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>
diff --git a/test/tint/bug/tint/1666.wgsl b/test/tint/bug/tint/1666.wgsl
new file mode 100644
index 0000000..ad56594
--- /dev/null
+++ b/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();
+}
diff --git a/test/tint/bug/tint/1666.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1666.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..3f8125c
--- /dev/null
+++ b/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};
+ ^
+
+
diff --git a/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..2a69fe5
--- /dev/null
+++ b/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
diff --git a/test/tint/bug/tint/1666.wgsl.expected.glsl b/test/tint/bug/tint/1666.wgsl.expected.glsl
new file mode 100644
index 0000000..2c6fa18
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/bug/tint/1666.wgsl.expected.msl b/test/tint/bug/tint/1666.wgsl.expected.msl
new file mode 100644
index 0000000..9327f8a
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/tint/bug/tint/1666.wgsl.expected.spvasm b/test/tint/bug/tint/1666.wgsl.expected.spvasm
new file mode 100644
index 0000000..ac0c05f
--- /dev/null
+++ b/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
diff --git a/test/tint/bug/tint/1666.wgsl.expected.wgsl b/test/tint/bug/tint/1666.wgsl.expected.wgsl
new file mode 100644
index 0000000..e6f5a2a
--- /dev/null
+++ b/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();
+}
