tint: Add test cases for tint:1666

This change does not attempt to fix this issue.

Bug: tint:1665
Bug: tint:1666
Change-Id: I9b40a25279b939977c826f38592518b6b086c06b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/101161
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>

test/tint/bug/tint/1666.wgsl

@@ -0,0 +1,29 @@
+fn vector() {
+    let idx = 3;
+    let x = vec2(1,2)[idx];
+}
+
+fn matrix() {
+    let idx = 4;
+    let x = mat2x2(1,2,3,4)[idx];
+}
+
+fn fixed_size_array() {
+    let arr = array(1,2);
+    let idx = 3;
+    let x = arr[idx];
+}
+
+@group(0) @binding(0) var<storage> rarr : array<f32>;
+fn runtime_size_array() {
+    let idx = -1;
+    let x = rarr[idx];
+}
+
+@compute @workgroup_size(1)
+fn f() {
+    vector();
+    matrix();
+    fixed_size_array();
+    runtime_size_array();
+}

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

@@ -0,0 +1,48 @@
+SKIP: FAILED
+
+void tint_symbol() {
+  const int idx = 3;
+  const int x = int2(1, 2)[idx];
+}
+
+void tint_symbol_1() {
+  const int idx = 4;
+  const float2 x = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[idx];
+}
+
+void fixed_size_array() {
+  const int arr[2] = {1, 2};
+  const int idx = 3;
+  const int x = arr[idx];
+}
+
+ByteAddressBuffer rarr : register(t0, space0);
+
+void runtime_size_array() {
+  const int idx = -1;
+  const float x = asfloat(rarr.Load((4u * uint(idx))));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  tint_symbol();
+  tint_symbol_1();
+  fixed_size_array();
+  runtime_size_array();
+  return;
+}
+DXC validation failure:
+shader.hlsl:3:28: error: vector element index '3' is out of bounds
+  const int x = int2(1, 2)[idx];
+                           ^
+shader.hlsl:8:69: error: matrix row index '4' is out of bounds
+  const float2 x = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[idx];
+                                                                    ^
+shader.hlsl:14:17: error: array index 3 is out of bounds
+  const int x = arr[idx];
+                ^
+shader.hlsl:12:3: note: array 'arr' declared here
+  const int arr[2] = {1, 2};
+  ^
+
+

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

@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void tint_symbol() {
+  const int idx = 3;
+  const int x = int2(1, 2)[idx];
+}
+
+void tint_symbol_1() {
+  const int idx = 4;
+  const float2 x = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[idx];
+}
+
+void fixed_size_array() {
+  const int arr[2] = {1, 2};
+  const int idx = 3;
+  const int x = arr[idx];
+}
+
+ByteAddressBuffer rarr : register(t0, space0);
+
+void runtime_size_array() {
+  const int idx = -1;
+  const float x = asfloat(rarr.Load((4u * uint(idx))));
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  tint_symbol();
+  tint_symbol_1();
+  fixed_size_array();
+  runtime_size_array();
+  return;
+}
+FXC validation failure:
+T:\tmp\dawn-temp\dawn-src\test\tint\Shader@0x0000019913EB4550(3,17-31): error X3030: array index out of bounds

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

@@ -0,0 +1,39 @@
+#version 310 es
+
+void vector() {
+  int idx = 3;
+  int x = ivec2(1, 2)[idx];
+}
+
+void matrix() {
+  int idx = 4;
+  vec2 x = mat2(vec2(1.0f, 2.0f), vec2(3.0f, 4.0f))[idx];
+}
+
+void fixed_size_array() {
+  int arr[2] = int[2](1, 2);
+  int idx = 3;
+  int x = arr[idx];
+}
+
+layout(binding = 0, std430) buffer rarr_block_ssbo {
+  float inner[];
+} rarr;
+
+void runtime_size_array() {
+  int idx = -1;
+  float x = rarr.inner[idx];
+}
+
+void f() {
+  vector();
+  matrix();
+  fixed_size_array();
+  runtime_size_array();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

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

@@ -0,0 +1,49 @@
+#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];
+};
+
+struct tint_symbol_3 {
+  /* 0x0000 */ tint_array<float, 1> arr;
+};
+
+void vector() {
+  int const idx = 3;
+  int const x = int2(1, 2)[idx];
+}
+
+void tint_symbol() {
+  int const idx = 4;
+  float2 const x = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[idx];
+}
+
+void fixed_size_array() {
+  tint_array<int, 2> const arr = tint_array<int, 2>{1, 2};
+  int const idx = 3;
+  int const x = arr[idx];
+}
+
+void runtime_size_array(const device tint_array<float, 1>* const tint_symbol_1) {
+  int const idx = -1;
+  float const x = (*(tint_symbol_1))[idx];
+}
+
+kernel void f(const device tint_symbol_3* tint_symbol_2 [[buffer(0)]]) {
+  vector();
+  tint_symbol();
+  fixed_size_array();
+  runtime_size_array(&((*(tint_symbol_2)).arr));
+  return;
+}
+

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

@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 60
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %rarr_block "rarr_block"
+               OpMemberName %rarr_block 0 "inner"
+               OpName %rarr "rarr"
+               OpName %vector "vector"
+               OpName %matrix "matrix"
+               OpName %var_for_index "var_for_index"
+               OpName %fixed_size_array "fixed_size_array"
+               OpName %var_for_index_1 "var_for_index_1"
+               OpName %runtime_size_array "runtime_size_array"
+               OpName %f "f"
+               OpDecorate %rarr_block Block
+               OpMemberDecorate %rarr_block 0 Offset 0
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpDecorate %rarr NonWritable
+               OpDecorate %rarr DescriptorSet 0
+               OpDecorate %rarr Binding 0
+               OpDecorate %_arr_int_uint_2 ArrayStride 4
+      %float = OpTypeFloat 32
+%_runtimearr_float = OpTypeRuntimeArray %float
+ %rarr_block = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_rarr_block = OpTypePointer StorageBuffer %rarr_block
+       %rarr = OpVariable %_ptr_StorageBuffer_rarr_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_3 = OpConstant %int 3
+      %v2int = OpTypeVector %int 2
+      %int_1 = OpConstant %int 1
+      %int_2 = OpConstant %int 2
+         %15 = OpConstantComposite %v2int %int_1 %int_2
+      %int_4 = OpConstant %int 4
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+    %float_1 = OpConstant %float 1
+    %float_2 = OpConstant %float 2
+         %24 = OpConstantComposite %v2float %float_1 %float_2
+    %float_3 = OpConstant %float 3
+    %float_4 = OpConstant %float 4
+         %27 = OpConstantComposite %v2float %float_3 %float_4
+         %28 = OpConstantComposite %mat2v2float %24 %27
+%_ptr_Function_mat2v2float = OpTypePointer Function %mat2v2float
+         %31 = OpConstantNull %mat2v2float
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+%_arr_int_uint_2 = OpTypeArray %int %uint_2
+         %40 = OpConstantComposite %_arr_int_uint_2 %int_1 %int_2
+%_ptr_Function__arr_int_uint_2 = OpTypePointer Function %_arr_int_uint_2
+         %43 = OpConstantNull %_arr_int_uint_2
+%_ptr_Function_int = OpTypePointer Function %int
+     %int_n1 = OpConstant %int -1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+     %vector = OpFunction %void None %6
+          %9 = OpLabel
+         %16 = OpVectorExtractDynamic %int %15 %int_3
+               OpReturn
+               OpFunctionEnd
+     %matrix = OpFunction %void None %6
+         %18 = OpLabel
+%var_for_index = OpVariable %_ptr_Function_mat2v2float Function %31
+               OpStore %var_for_index %28
+         %33 = OpAccessChain %_ptr_Function_v2float %var_for_index %int_4
+         %34 = OpLoad %v2float %33
+               OpReturn
+               OpFunctionEnd
+%fixed_size_array = OpFunction %void None %6
+         %36 = OpLabel
+%var_for_index_1 = OpVariable %_ptr_Function__arr_int_uint_2 Function %43
+               OpStore %var_for_index_1 %40
+         %45 = OpAccessChain %_ptr_Function_int %var_for_index_1 %int_3
+         %46 = OpLoad %int %45
+               OpReturn
+               OpFunctionEnd
+%runtime_size_array = OpFunction %void None %6
+         %48 = OpLabel
+         %52 = OpAccessChain %_ptr_StorageBuffer_float %rarr %uint_0 %int_n1
+         %53 = OpLoad %float %52
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %6
+         %55 = OpLabel
+         %56 = OpFunctionCall %void %vector
+         %57 = OpFunctionCall %void %matrix
+         %58 = OpFunctionCall %void %fixed_size_array
+         %59 = OpFunctionCall %void %runtime_size_array
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1666.wgsl.expected.wgsl

@@ -0,0 +1,30 @@
+fn vector() {
+  let idx = 3;
+  let x = vec2(1, 2)[idx];
+}
+
+fn matrix() {
+  let idx = 4;
+  let x = mat2x2(1, 2, 3, 4)[idx];
+}
+
+fn fixed_size_array() {
+  let arr = array(1, 2);
+  let idx = 3;
+  let x = arr[idx];
+}
+
+@group(0) @binding(0) var<storage> rarr : array<f32>;
+
+fn runtime_size_array() {
+  let idx = -1;
+  let x = rarr[idx];
+}
+
+@compute @workgroup_size(1)
+fn f() {
+  vector();
+  matrix();
+  fixed_size_array();
+  runtime_size_array();
+}