test/tint: Add tests showing tint:1474 is fixed

Fixed: tint:1474
Change-Id: I462ae7bdb17a394b0a4c185fc815fe9727044555
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/105963
Commit-Queue: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>

test/tint/bug/tint/1474-a.wgsl

@@ -0,0 +1,13 @@
+// Example taken from https://github.com/gpuweb/gpuweb/pull/2622
+@compute @workgroup_size(1,1,1)
+fn main() {
+    const cond = true;
+    while (cond) {
+        if cond {
+            break;
+        } else {
+            return;
+        } // Overall behavior is {Break, Return}
+    }
+    let x = 5; // valid, but behavior is {Break, Return}
+}

test/tint/bug/tint/1474-a.wgsl.expected.dxc.hlsl

@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void main() {
+  while(true) {
+    if (true) {
+      break;
+    } else {
+      return;
+    }
+  }
+  const int x = 5;
+  return;
+}

test/tint/bug/tint/1474-a.wgsl.expected.fxc.hlsl

@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void main() {
+  while(true) {
+    if (true) {
+      break;
+    } else {
+      return;
+    }
+  }
+  const int x = 5;
+  return;
+}

test/tint/bug/tint/1474-a.wgsl.expected.glsl

@@ -0,0 +1,18 @@
+#version 310 es
+
+void tint_symbol() {
+  while(true) {
+    if (true) {
+      break;
+    } else {
+      return;
+    }
+  }
+  int x = 5;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/bug/tint/1474-a.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  while(true) {
+    if (true) {
+      break;
+    } else {
+      return;
+    }
+  }
+  int const x = 5;
+  return;
+}
+

test/tint/bug/tint/1474-a.wgsl.expected.spvasm

@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %bool = OpTypeBool
+         %10 = OpConstantNull %bool
+       %true = OpConstantTrue %bool
+        %int = OpTypeInt 32 1
+      %int_5 = OpConstant %int 5
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+               OpBranch %5
+          %5 = OpLabel
+               OpLoopMerge %6 %7 None
+               OpBranch %8
+          %8 = OpLabel
+               OpSelectionMerge %11 None
+               OpBranchConditional %10 %12 %11
+         %12 = OpLabel
+               OpBranch %6
+         %11 = OpLabel
+               OpSelectionMerge %14 None
+               OpBranchConditional %true %15 %16
+         %15 = OpLabel
+               OpBranch %6
+         %16 = OpLabel
+               OpReturn
+         %14 = OpLabel
+               OpBranch %7
+          %7 = OpLabel
+               OpBranch %5
+          %6 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1474-a.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+@compute @workgroup_size(1, 1, 1)
+fn main() {
+  const cond = true;
+  while(cond) {
+    if (cond) {
+      break;
+    } else {
+      return;
+    }
+  }
+  let x = 5;
+}

test/tint/bug/tint/1474-b.wgsl

@@ -0,0 +1,12 @@
+// Example taken from https://github.com/gpuweb/gpuweb/pull/2622
+@group(0) @binding(0) var<storage, read_write> non_uniform_value : i32;
+
+@compute @workgroup_size(1,1,1)
+fn main() {
+    return;
+    let non_uniform_cond = non_uniform_value == 0;
+    if (non_uniform_cond) {
+        workgroupBarrier(); // valid, unreachable code does not contribute
+                            // to the uniformity analysis
+    }
+}

test/tint/bug/tint/1474-b.wgsl.expected.dxc.hlsl

@@ -0,0 +1,15 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+RWByteAddressBuffer non_uniform_value : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  return;
+  const bool non_uniform_cond = (asint(non_uniform_value.Load(0u)) == 0);
+  if (non_uniform_cond) {
+    GroupMemoryBarrierWithGroupSync();
+  }
+  return;
+}

test/tint/bug/tint/1474-b.wgsl.expected.fxc.hlsl

@@ -0,0 +1,15 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+RWByteAddressBuffer non_uniform_value : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  return;
+  const bool non_uniform_cond = (asint(non_uniform_value.Load(0u)) == 0);
+  if (non_uniform_cond) {
+    GroupMemoryBarrierWithGroupSync();
+  }
+  return;
+}

test/tint/bug/tint/1474-b.wgsl.expected.glsl

@@ -0,0 +1,23 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+#version 310 es
+
+layout(binding = 0, std430) buffer non_uniform_value_block_ssbo {
+  int inner;
+} non_uniform_value;
+
+void tint_symbol() {
+  return;
+  bool non_uniform_cond = (non_uniform_value.inner == 0);
+  if (non_uniform_cond) {
+    barrier();
+  }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/bug/tint/1474-b.wgsl.expected.msl

@@ -0,0 +1,16 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(0)]]) {
+  return;
+  bool const non_uniform_cond = (*(tint_symbol_1) == 0);
+  if (non_uniform_cond) {
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+  }
+  return;
+}
+

test/tint/bug/tint/1474-b.wgsl.expected.spvasm

@@ -0,0 +1,31 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %non_uniform_value_block "non_uniform_value_block"
+               OpMemberName %non_uniform_value_block 0 "inner"
+               OpName %non_uniform_value "non_uniform_value"
+               OpName %main "main"
+               OpDecorate %non_uniform_value_block Block
+               OpMemberDecorate %non_uniform_value_block 0 Offset 0
+               OpDecorate %non_uniform_value DescriptorSet 0
+               OpDecorate %non_uniform_value Binding 0
+        %int = OpTypeInt 32 1
+%non_uniform_value_block = OpTypeStruct %int
+%_ptr_StorageBuffer_non_uniform_value_block = OpTypePointer StorageBuffer %non_uniform_value_block
+%non_uniform_value = OpVariable %_ptr_StorageBuffer_non_uniform_value_block StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1474-b.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+bug/tint/1474-b.wgsl:7:9 warning: code is unreachable
+    let non_uniform_cond = non_uniform_value == 0;
+        ^^^^^^^^^^^^^^^^
+
+@group(0) @binding(0) var<storage, read_write> non_uniform_value : i32;
+
+@compute @workgroup_size(1, 1, 1)
+fn main() {
+  return;
+  let non_uniform_cond = (non_uniform_value == 0);
+  if (non_uniform_cond) {
+    workgroupBarrier();
+  }
+}