tint: add e2e test for crbug.com/tint/1605

Bug: tint:1605
Change-Id: I6a94552110f4864dcd7fe011063fa35606104b48
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/103984
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

test/tint/bug/tint/1605.wgsl

@@ -0,0 +1,16 @@
+@group(0) @binding(0)
+var<uniform> b: i32;
+
+fn func_3() -> bool {
+    for (var i = 0; i < b; i++) {
+        for (var j = -1; j == 1; j++) {
+            return false;
+        }
+    }
+    return false;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+    func_3();
+}

test/tint/bug/tint/1605.wgsl.expected.dxc.hlsl

@@ -0,0 +1,22 @@
+cbuffer cbuffer_b : register(b0, space0) {
+  uint4 b[1];
+};
+
+bool func_3() {
+  {
+    [loop] for(int i = 0; (i < asint(b[0].x)); i = (i + 1)) {
+      {
+        [loop] for(int j = -1; (j == 1); j = (j + 1)) {
+          return false;
+        }
+      }
+    }
+  }
+  return false;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  func_3();
+  return;
+}

test/tint/bug/tint/1605.wgsl.expected.fxc.hlsl

@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+cbuffer cbuffer_b : register(b0, space0) {
+  uint4 b[1];
+};
+
+bool func_3() {
+  {
+    [loop] for(int i = 0; (i < asint(b[0].x)); i = (i + 1)) {
+      {
+        [loop] for(int j = -1; (j == 1); j = (j + 1)) {
+          return false;
+        }
+      }
+    }
+  }
+  return false;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  func_3();
+  return;
+}
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3557: loop doesn't seem to do anything, consider removing [loop]
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3551: infinite loop detected - loop writes no values
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3557: loop doesn't seem to do anything, consider removing [loop]
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3551: infinite loop detected - loop writes no values
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3557: loop doesn't seem to do anything, consider removing [loop]
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3551: infinite loop detected - loop writes no values
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3557: loop only executes for 0 iteration(s), consider removing [loop]
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(9,16-53): warning X3557: loop doesn't seem to do anything, consider removing [loop]
+C:\src\dawn\test\tint\Shader@0x0000026A74A16540(5,11-13): error X4555: cannot use casts on l-values
+

test/tint/bug/tint/1605.wgsl.expected.glsl

@@ -0,0 +1,28 @@
+#version 310 es
+
+layout(binding = 0, std140) uniform b_block_ubo {
+  int inner;
+} b;
+
+bool func_3() {
+  {
+    for(int i = 0; (i < b.inner); i = (i + 1)) {
+      {
+        for(int j = -1; (j == 1); j = (j + 1)) {
+          return false;
+        }
+      }
+    }
+  }
+  return false;
+}
+
+void tint_symbol() {
+  func_3();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

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

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+bool func_3(const constant int* const tint_symbol_1) {
+  for(int i = 0; (i < *(tint_symbol_1)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    for(int j = -1; (j == 1); j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
+      return false;
+    }
+  }
+  return false;
+}
+
+kernel void tint_symbol(const constant int* tint_symbol_2 [[buffer(0)]]) {
+  func_3(tint_symbol_2);
+  return;
+}
+

test/tint/bug/tint/1605.wgsl.expected.spvasm

@@ -0,0 +1,92 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %b_block "b_block"
+               OpMemberName %b_block 0 "inner"
+               OpName %b "b"
+               OpName %func_3 "func_3"
+               OpName %i "i"
+               OpName %j "j"
+               OpName %main "main"
+               OpDecorate %b_block Block
+               OpMemberDecorate %b_block 0 Offset 0
+               OpDecorate %b NonWritable
+               OpDecorate %b DescriptorSet 0
+               OpDecorate %b Binding 0
+        %int = OpTypeInt 32 1
+    %b_block = OpTypeStruct %int
+%_ptr_Uniform_b_block = OpTypePointer Uniform %b_block
+          %b = OpVariable %_ptr_Uniform_b_block Uniform
+       %bool = OpTypeBool
+          %5 = OpTypeFunction %bool
+          %9 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %int_n1 = OpConstant %int -1
+      %int_1 = OpConstant %int 1
+         %38 = OpConstantNull %bool
+       %void = OpTypeVoid
+         %43 = OpTypeFunction %void
+     %func_3 = OpFunction %bool None %5
+          %8 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %9
+          %j = OpVariable %_ptr_Function_int Function %9
+               OpStore %i %9
+               OpBranch %12
+         %12 = OpLabel
+               OpLoopMerge %13 %14 None
+               OpBranch %15
+         %15 = OpLabel
+         %17 = OpLoad %int %i
+         %21 = OpAccessChain %_ptr_Uniform_int %b %uint_0
+         %22 = OpLoad %int %21
+         %23 = OpSLessThan %bool %17 %22
+         %16 = OpLogicalNot %bool %23
+               OpSelectionMerge %24 None
+               OpBranchConditional %16 %25 %24
+         %25 = OpLabel
+               OpBranch %13
+         %24 = OpLabel
+               OpStore %j %int_n1
+               OpBranch %28
+         %28 = OpLabel
+               OpLoopMerge %29 %30 None
+               OpBranch %31
+         %31 = OpLabel
+         %33 = OpLoad %int %j
+         %35 = OpIEqual %bool %33 %int_1
+         %32 = OpLogicalNot %bool %35
+               OpSelectionMerge %36 None
+               OpBranchConditional %32 %37 %36
+         %37 = OpLabel
+               OpBranch %29
+         %36 = OpLabel
+               OpReturnValue %38
+         %30 = OpLabel
+         %39 = OpLoad %int %j
+         %40 = OpIAdd %int %39 %int_1
+               OpStore %j %40
+               OpBranch %28
+         %29 = OpLabel
+               OpBranch %14
+         %14 = OpLabel
+         %41 = OpLoad %int %i
+         %42 = OpIAdd %int %41 %int_1
+               OpStore %i %42
+               OpBranch %12
+         %13 = OpLabel
+               OpReturnValue %38
+               OpFunctionEnd
+       %main = OpFunction %void None %43
+         %46 = OpLabel
+         %47 = OpFunctionCall %bool %func_3
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1605.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+@group(0) @binding(0) var<uniform> b : i32;
+
+fn func_3() -> bool {
+  for(var i = 0; (i < b); i++) {
+    for(var j = -1; (j == 1); j++) {
+      return false;
+    }
+  }
+  return false;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  func_3();
+}