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 <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/tint/bug/tint/1604.wgsl

@@ -0,0 +1,15 @@
+@group(0) @binding(0)
+var<uniform> x: i32;
+
+@compute @workgroup_size(1)
+fn main() {
+    switch (x) {
+        case 0: {
+            loop {
+                return;
+            }
+        }
+        default: {
+        }
+    }
+}

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;
+}

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.
+

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;
+}

test/tint/bug/tint/1604.wgsl.expected.msl

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+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;
+}
+

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

test/tint/bug/tint/1604.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+@group(0) @binding(0) var<uniform> x : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+  switch(x) {
+    case 0: {
+      loop {
+        return;
+      }
+    }
+    default: {
+    }
+  }
+}