From f2c1d0aa5bd8d91d49ee7e2726186a66500f2a39 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 21 Sep 2022 17:19:04 +0000 Subject: [PATCH] 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 Kokoro: Kokoro Reviewed-by: Antonio Maiorano --- test/tint/bug/tint/1666.wgsl | 29 ++++++ .../tint/bug/tint/1666.wgsl.expected.dxc.hlsl | 48 +++++++++ .../tint/bug/tint/1666.wgsl.expected.fxc.hlsl | 35 +++++++ test/tint/bug/tint/1666.wgsl.expected.glsl | 39 ++++++++ test/tint/bug/tint/1666.wgsl.expected.msl | 49 ++++++++++ test/tint/bug/tint/1666.wgsl.expected.spvasm | 97 +++++++++++++++++++ test/tint/bug/tint/1666.wgsl.expected.wgsl | 30 ++++++ 7 files changed, 327 insertions(+) create mode 100644 test/tint/bug/tint/1666.wgsl create mode 100644 test/tint/bug/tint/1666.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1666.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1666.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1666.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1666.wgsl.expected.wgsl diff --git a/test/tint/bug/tint/1666.wgsl b/test/tint/bug/tint/1666.wgsl new file mode 100644 index 0000000000..ad565944e8 --- /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 rarr : array; +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 0000000000..3f8125cbf8 --- /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 0000000000..2a69fe57af --- /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 0000000000..2c6fa18315 --- /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 0000000000..9327f8a856 --- /dev/null +++ b/test/tint/bug/tint/1666.wgsl.expected.msl @@ -0,0 +1,49 @@ +#include + +using namespace metal; + +template +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 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 const arr = tint_array{1, 2}; + int const idx = 3; + int const x = arr[idx]; +} + +void runtime_size_array(const device tint_array* 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 0000000000..ac0c05f241 --- /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 0000000000..e6f5a2a01e --- /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 rarr : array; + +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(); +}