diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 122f50e414..c66278f46f 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -344,6 +344,181 @@ bool GeneratorImpl::EmitDynamicVectorAssignment( return true; } +bool GeneratorImpl::EmitDynamicMatrixVectorAssignment( + const ast::AssignmentStatement* stmt, + const sem::Matrix* mat) { + auto name = utils::GetOrCreate( + dynamic_matrix_vector_write_, mat, [&]() -> std::string { + std::string fn; + { + std::ostringstream ss; + if (!EmitType(ss, mat, tint::ast::StorageClass::kInvalid, + ast::Access::kUndefined, "")) { + return ""; + } + fn = UniqueIdentifier("set_vector_" + ss.str()); + } + { + auto out = line(&helpers_); + out << "void " << fn << "(inout "; + if (!EmitTypeAndName(out, mat, ast::StorageClass::kInvalid, + ast::Access::kUndefined, "mat")) { + return ""; + } + out << ", int col, "; + if (!EmitTypeAndName(out, mat->ColumnType(), + ast::StorageClass::kInvalid, + ast::Access::kUndefined, "val")) { + return ""; + } + out << ") {"; + } + { + ScopedIndent si(&helpers_); + line(&helpers_) << "switch (col) {"; + { + ScopedIndent si2(&helpers_); + for (uint32_t i = 0; i < mat->columns(); ++i) { + line(&helpers_) + << "case " << i << ": mat[" << i << "] = val; break;"; + } + } + line(&helpers_) << "}"; + } + line(&helpers_) << "}"; + line(&helpers_); + return fn; + }); + + if (name.empty()) { + return false; + } + + auto* ast_access_expr = stmt->lhs->As(); + + auto out = line(); + out << name << "("; + if (!EmitExpression(out, ast_access_expr->object)) { + return false; + } + out << ", "; + if (!EmitExpression(out, ast_access_expr->index)) { + return false; + } + out << ", "; + if (!EmitExpression(out, stmt->rhs)) { + return false; + } + out << ");"; + + return true; +} + +bool GeneratorImpl::EmitDynamicMatrixScalarAssignment( + const ast::AssignmentStatement* stmt, + const sem::Matrix* mat) { + auto* lhs_col_access = stmt->lhs->As(); + auto* lhs_row_access = + lhs_col_access->object->As(); + + auto name = utils::GetOrCreate( + dynamic_matrix_scalar_write_, mat, [&]() -> std::string { + std::string fn; + { + std::ostringstream ss; + if (!EmitType(ss, mat, tint::ast::StorageClass::kInvalid, + ast::Access::kUndefined, "")) { + return ""; + } + fn = UniqueIdentifier("set_scalar_" + ss.str()); + } + { + auto out = line(&helpers_); + out << "void " << fn << "(inout "; + if (!EmitTypeAndName(out, mat, ast::StorageClass::kInvalid, + ast::Access::kUndefined, "mat")) { + return ""; + } + out << ", int col, int row, "; + if (!EmitTypeAndName(out, mat->type(), ast::StorageClass::kInvalid, + ast::Access::kUndefined, "val")) { + return ""; + } + out << ") {"; + } + { + ScopedIndent si(&helpers_); + line(&helpers_) << "switch (col) {"; + { + ScopedIndent si2(&helpers_); + auto* vec = + TypeOf(lhs_row_access->object)->UnwrapRef()->As(); + for (uint32_t i = 0; i < mat->columns(); ++i) { + line(&helpers_) << "case " << i << ":"; + { + auto vec_name = "mat[" + std::to_string(i) + "]"; + ScopedIndent si3(&helpers_); + { + auto out = line(&helpers_); + switch (mat->rows()) { + case 2: + out << vec_name + << " = (row.xx == int2(0, 1)) ? val.xx : " << vec_name + << ";"; + break; + case 3: + out << vec_name + << " = (row.xxx == int3(0, 1, 2)) ? val.xxx : " + << vec_name << ";"; + break; + case 4: + out << vec_name + << " = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : " + << vec_name << ";"; + break; + default: + TINT_UNREACHABLE(Writer, builder_.Diagnostics()) + << "invalid vector size " << vec->Width(); + break; + } + } + line(&helpers_) << "break;"; + } + } + } + line(&helpers_) << "}"; + } + line(&helpers_) << "}"; + line(&helpers_); + return fn; + }); + + if (name.empty()) { + return false; + } + + auto out = line(); + out << name << "("; + if (!EmitExpression(out, lhs_row_access->object)) { + return false; + } + out << ", "; + if (!EmitExpression(out, lhs_col_access->index)) { + return false; + } + out << ", "; + if (!EmitExpression(out, lhs_row_access->index)) { + return false; + } + out << ", "; + if (!EmitExpression(out, stmt->rhs)) { + return false; + } + out << ");"; + + return true; +} + bool GeneratorImpl::EmitIndexAccessor( std::ostream& out, const ast::IndexAccessorExpression* expr) { @@ -387,9 +562,34 @@ bool GeneratorImpl::EmitBitcast(std::ostream& out, } bool GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) { - if (auto* idx = stmt->lhs->As()) { - if (auto* vec = TypeOf(idx->object)->UnwrapRef()->As()) { - auto* rhs_sem = builder_.Sem().Get(idx->index); + if (auto* lhs_access = stmt->lhs->As()) { + // BUG(crbug.com/tint/1333): work around assignment of scalar to matrices + // with at least one dynamic index + if (auto* lhs_sub_access = + lhs_access->object->As()) { + if (auto* mat = + TypeOf(lhs_sub_access->object)->UnwrapRef()->As()) { + auto* rhs_col_idx_sem = builder_.Sem().Get(lhs_access->index); + auto* rhs_row_idx_sem = builder_.Sem().Get(lhs_sub_access->index); + if (!rhs_col_idx_sem->ConstantValue().IsValid() || + !rhs_row_idx_sem->ConstantValue().IsValid()) { + return EmitDynamicMatrixScalarAssignment(stmt, mat); + } + } + } + // BUG(crbug.com/tint/1333): work around assignment of vector to matrices + // with dynamic indices + const auto* lhs_access_type = TypeOf(lhs_access->object)->UnwrapRef(); + if (auto* mat = lhs_access_type->As()) { + auto* lhs_index_sem = builder_.Sem().Get(lhs_access->index); + if (!lhs_index_sem->ConstantValue().IsValid()) { + return EmitDynamicMatrixVectorAssignment(stmt, mat); + } + } + // BUG(crbug.com/tint/534): work around assignment to vectors with dynamic + // indices + if (auto* vec = lhs_access_type->As()) { + auto* rhs_sem = builder_.Sem().Get(lhs_access->index); if (!rhs_sem->ConstantValue().IsValid()) { return EmitDynamicVectorAssignment(stmt, vec); } diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index c2532a2dec..12124d2ecf 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -423,6 +423,26 @@ class GeneratorImpl : public TextGenerator { /// @returns true on success bool EmitDynamicVectorAssignment(const ast::AssignmentStatement* stmt, const sem::Vector* vec); + /// Emits call to a helper matrix assignment function for the input assignment + /// statement and matrix type. This is used to work around FXC issues where + /// assignment of a vector to a matrix with a dynamic index causes compilation + /// failures. + /// @param stmt assignment statement that corresponds to a matrix assignment + /// via an accessor expression + /// @param mat the matrix type being assigned to + /// @returns true on success + bool EmitDynamicMatrixVectorAssignment(const ast::AssignmentStatement* stmt, + const sem::Matrix* mat); + /// Emits call to a helper matrix assignment function for the input assignment + /// statement and matrix type. This is used to work around FXC issues where + /// assignment of a scalar to a matrix with at least one dynamic index causes + /// compilation failures. + /// @param stmt assignment statement that corresponds to a matrix assignment + /// via an accessor expression + /// @param mat the matrix type being assigned to + /// @returns true on success + bool EmitDynamicMatrixScalarAssignment(const ast::AssignmentStatement* stmt, + const sem::Matrix* mat); /// Handles generating a builtin method name /// @param intrinsic the semantic info for the intrinsic @@ -491,6 +511,10 @@ class GeneratorImpl : public TextGenerator { std::unordered_map intrinsics_; std::unordered_map structure_builders_; std::unordered_map dynamic_vector_write_; + std::unordered_map + dynamic_matrix_vector_write_; + std::unordered_map + dynamic_matrix_scalar_write_; }; } // namespace hlsl diff --git a/src/writer/hlsl/generator_impl_assign_test.cc b/src/writer/hlsl/generator_impl_assign_test.cc index f5a0711524..98d9eb6b5c 100644 --- a/src/writer/hlsl/generator_impl_assign_test.cc +++ b/src/writer/hlsl/generator_impl_assign_test.cc @@ -22,17 +22,187 @@ namespace { using HlslGeneratorImplTest_Assign = TestHelper; TEST_F(HlslGeneratorImplTest_Assign, Emit_Assign) { - Global("lhs", ty.i32(), ast::StorageClass::kPrivate); - Global("rhs", ty.i32(), ast::StorageClass::kPrivate); - auto* assign = Assign("lhs", "rhs"); - WrapInFunction(assign); + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.i32())), + Decl(Var("rhs", ty.i32())), + Assign("lhs", "rhs"), + }); GeneratorImpl& gen = Build(); - gen.increment_indent(); + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ(gen.result(), + R"(void fn() { + int lhs = 0; + int rhs = 0; + lhs = rhs; +} +)"); +} - ASSERT_TRUE(gen.EmitStatement(assign)) << gen.error(); - EXPECT_EQ(gen.result(), " lhs = rhs;\n"); +TEST_F(HlslGeneratorImplTest_Assign, Emit_Vector_Assign_ConstantIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.vec3())), + Decl(Var("rhs", ty.f32())), + Decl(Const("index", ty.u32(), Expr(0u))), + Assign(IndexAccessor("lhs", "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ(gen.result(), + R"(void fn() { + float3 lhs = float3(0.0f, 0.0f, 0.0f); + float rhs = 0.0f; + const uint index = 0u; + lhs[index] = rhs; +} +)"); +} + +TEST_F(HlslGeneratorImplTest_Assign, Emit_Vector_Assign_DynamicIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.vec3())), + Decl(Var("rhs", ty.f32())), + Decl(Var("index", ty.u32())), + Assign(IndexAccessor("lhs", "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ(gen.result(), + R"(void set_float3(inout float3 vec, int idx, float val) { + vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; +} + +void fn() { + float3 lhs = float3(0.0f, 0.0f, 0.0f); + float rhs = 0.0f; + uint index = 0u; + set_float3(lhs, index, rhs); +} +)"); +} + +TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Vector_ConstantIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.mat4x2())), + Decl(Var("rhs", ty.vec2())), + Decl(Const("index", ty.u32(), Expr(0u))), + Assign(IndexAccessor("lhs", "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ(gen.result(), + R"(void fn() { + float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + float2 rhs = float2(0.0f, 0.0f); + const uint index = 0u; + lhs[index] = rhs; +} +)"); +} + +TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Vector_DynamicIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.mat4x2())), + Decl(Var("rhs", ty.vec2())), + Decl(Var("index", ty.u32())), + Assign(IndexAccessor("lhs", "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ( + gen.result(), + R"(void set_vector_float4x2(inout float4x2 mat, int col, float2 val) { + switch (col) { + case 0: mat[0] = val; break; + case 1: mat[1] = val; break; + case 2: mat[2] = val; break; + case 3: mat[3] = val; break; + } +} + +void fn() { + float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + float2 rhs = float2(0.0f, 0.0f); + uint index = 0u; + set_vector_float4x2(lhs, index, rhs); +} +)"); +} + +TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Scalar_ConstantIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.mat4x2())), + Decl(Var("rhs", ty.f32())), + Decl(Const("index", ty.u32(), Expr(0u))), + Assign(IndexAccessor(IndexAccessor("lhs", "index"), "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ(gen.result(), + R"(void fn() { + float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + float rhs = 0.0f; + const uint index = 0u; + lhs[index][index] = rhs; +} +)"); +} + +TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Scalar_DynamicIndex) { + Func("fn", {}, ty.void_(), + { + Decl(Var("lhs", ty.mat4x2())), + Decl(Var("rhs", ty.f32())), + Decl(Var("index", ty.u32())), + Assign(IndexAccessor(IndexAccessor("lhs", "index"), "index"), "rhs"), + }); + + GeneratorImpl& gen = Build(); + + ASSERT_TRUE(gen.Generate()); + EXPECT_EQ( + gen.result(), + R"(void set_scalar_float4x2(inout float4x2 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xx == int2(0, 1)) ? val.xx : mat[0]; + break; + case 1: + mat[1] = (row.xx == int2(0, 1)) ? val.xx : mat[1]; + break; + case 2: + mat[2] = (row.xx == int2(0, 1)) ? val.xx : mat[2]; + break; + case 3: + mat[3] = (row.xx == int2(0, 1)) ? val.xx : mat[3]; + break; + } +} + +void fn() { + float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + float rhs = 0.0f; + uint index = 0u; + set_scalar_float4x2(lhs, index, index, rhs); +} +)"); } } // namespace diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl new file mode 100644 index 0000000000..7c39426c43 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl @@ -0,0 +1,12 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i][0] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.hlsl new file mode 100644 index 0000000000..d9c9833994 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + set_scalar_float2x4(m1, 0, uniforms[0].x, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.msl new file mode 100644 index 0000000000..82056b44cf --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + float2x4 m1 = float2x4(0.0f); + m1[(*(tint_symbol_1)).i][0] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.spvasm new file mode 100644 index 0000000000..c2b242a8ee --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.spvasm @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %m1 "m1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float + %14 = OpConstantNull %mat2v4float + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %m1 = OpVariable %_ptr_Function_mat2v4float Function %14 + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %22 = OpAccessChain %_ptr_Function_float %m1 %18 %int_0 + OpStore %22 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.wgsl new file mode 100644 index 0000000000..eeaa5a97d0 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.wgsl @@ -0,0 +1,13 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i][0] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl new file mode 100644 index 0000000000..9f5e96196e --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl @@ -0,0 +1,12 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.hlsl new file mode 100644 index 0000000000..254519f61e --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + set_scalar_float2x4(m1, uniforms[0].y, uniforms[0].x, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.msl new file mode 100644 index 0000000000..ea2b9bae7e --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + float2x4 m1 = float2x4(0.0f); + m1[(*(tint_symbol_1)).i][(*(tint_symbol_1)).j] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.spvasm new file mode 100644 index 0000000000..b3c73b46b7 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.spvasm @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %m1 "m1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float + %14 = OpConstantNull %mat2v4float + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %m1 = OpVariable %_ptr_Function_mat2v4float Function %14 + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %21 = OpLoad %uint %20 + %23 = OpAccessChain %_ptr_Function_float %m1 %18 %21 + OpStore %23 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.wgsl new file mode 100644 index 0000000000..b4c8fe9726 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.wgsl @@ -0,0 +1,13 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl new file mode 100644 index 0000000000..345e6e065d --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl @@ -0,0 +1,13 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[0][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.hlsl new file mode 100644 index 0000000000..75a62179b4 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; +static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + +[numthreads(1, 1, 1)] +void main() { + set_scalar_float2x4(m1, uniforms[0].y, 0, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.msl new file mode 100644 index 0000000000..4e4bb57d55 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + thread float2x4 tint_symbol_1 = float2x4(0.0f); + tint_symbol_1[0][(*(tint_symbol_2)).j] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.spvasm new file mode 100644 index 0000000000..3f68ce15e1 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.spvasm @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %m1 "m1" + OpName %main "main" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float + %10 = OpConstantNull %mat2v4float + %m1 = OpVariable %_ptr_Private_mat2v4float Private %10 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %11 + %14 = OpLabel + %19 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %20 = OpLoad %uint %19 + %22 = OpAccessChain %_ptr_Private_float %m1 %int_0 %20 + OpStore %22 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.wgsl new file mode 100644 index 0000000000..9b78e5a4a5 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[0][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl new file mode 100644 index 0000000000..fa0ad2531d --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl @@ -0,0 +1,12 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i] = vec4(1.0); +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.hlsl new file mode 100644 index 0000000000..4b5af7190b --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.hlsl @@ -0,0 +1,17 @@ +void set_vector_float2x4(inout float2x4 mat, int col, float4 val) { + switch (col) { + case 0: mat[0] = val; break; + case 1: mat[1] = val; break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + set_vector_float2x4(m1, uniforms[0].x, float4((1.0f).xxxx)); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.msl new file mode 100644 index 0000000000..a5aed7f46e --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + float2x4 m1 = float2x4(0.0f); + m1[(*(tint_symbol_1)).i] = float4(1.0f); + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.spvasm new file mode 100644 index 0000000000..afc1c44a7e --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.spvasm @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %m1 "m1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float + %14 = OpConstantNull %mat2v4float + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Function_v4float = OpTypePointer Function %v4float + %float_1 = OpConstant %float 1 + %22 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1 + %main = OpFunction %void None %5 + %8 = OpLabel + %m1 = OpVariable %_ptr_Function_mat2v4float Function %14 + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %20 = OpAccessChain %_ptr_Function_v4float %m1 %18 + OpStore %20 %22 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.wgsl new file mode 100644 index 0000000000..ef9e7ad103 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.wgsl @@ -0,0 +1,13 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var m1 : mat2x4; + m1[uniforms.i] = vec4(1.0); +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl new file mode 100644 index 0000000000..53253314fa --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl @@ -0,0 +1,13 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i][0] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.hlsl new file mode 100644 index 0000000000..c52f157379 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; +static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + +[numthreads(1, 1, 1)] +void main() { + set_scalar_float2x4(m1, 0, uniforms[0].x, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.msl new file mode 100644 index 0000000000..31c3003e62 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + thread float2x4 tint_symbol_1 = float2x4(0.0f); + tint_symbol_1[(*(tint_symbol_2)).i][0] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.spvasm new file mode 100644 index 0000000000..b2bb365694 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.spvasm @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %m1 "m1" + OpName %main "main" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float + %10 = OpConstantNull %mat2v4float + %m1 = OpVariable %_ptr_Private_mat2v4float Private %10 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %11 + %14 = OpLabel + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %22 = OpAccessChain %_ptr_Private_float %m1 %18 %int_0 + OpStore %22 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.wgsl new file mode 100644 index 0000000000..ea6fc1b2ce --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i][0] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl new file mode 100644 index 0000000000..9ccb54a1df --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl @@ -0,0 +1,13 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.hlsl new file mode 100644 index 0000000000..1487ddd1d9 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; +static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + +[numthreads(1, 1, 1)] +void main() { + set_scalar_float2x4(m1, uniforms[0].y, uniforms[0].x, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.msl new file mode 100644 index 0000000000..7215cb0002 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + thread float2x4 tint_symbol_1 = float2x4(0.0f); + tint_symbol_1[(*(tint_symbol_2)).i][(*(tint_symbol_2)).j] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.spvasm new file mode 100644 index 0000000000..1b93b116f7 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.spvasm @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %m1 "m1" + OpName %main "main" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float + %10 = OpConstantNull %mat2v4float + %m1 = OpVariable %_ptr_Private_mat2v4float Private %10 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %11 + %14 = OpLabel + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %21 = OpLoad %uint %20 + %23 = OpAccessChain %_ptr_Private_float %m1 %18 %21 + OpStore %23 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.wgsl new file mode 100644 index 0000000000..b9d60a2bad --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl new file mode 100644 index 0000000000..345e6e065d --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl @@ -0,0 +1,13 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[0][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.hlsl new file mode 100644 index 0000000000..75a62179b4 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0]; + break; + case 1: + mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1]; + break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; +static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + +[numthreads(1, 1, 1)] +void main() { + set_scalar_float2x4(m1, uniforms[0].y, 0, 1.0f); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.msl new file mode 100644 index 0000000000..4e4bb57d55 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + thread float2x4 tint_symbol_1 = float2x4(0.0f); + tint_symbol_1[0][(*(tint_symbol_2)).j] = 1.0f; + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.spvasm new file mode 100644 index 0000000000..3f68ce15e1 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.spvasm @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %m1 "m1" + OpName %main "main" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float + %10 = OpConstantNull %mat2v4float + %m1 = OpVariable %_ptr_Private_mat2v4float Private %10 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Private_float = OpTypePointer Private %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %11 + %14 = OpLabel + %19 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %20 = OpLoad %uint %19 + %22 = OpAccessChain %_ptr_Private_float %m1 %int_0 %20 + OpStore %22 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.wgsl new file mode 100644 index 0000000000..9b78e5a4a5 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[0][uniforms.j] = 1.0; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl new file mode 100644 index 0000000000..d5b3c63fa8 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl @@ -0,0 +1,13 @@ +[[block]] struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i] = vec4(1.0); +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.hlsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.hlsl new file mode 100644 index 0000000000..bf033730c7 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.hlsl @@ -0,0 +1,17 @@ +void set_vector_float2x4(inout float2x4 mat, int col, float4 val) { + switch (col) { + case 0: mat[0] = val; break; + case 1: mat[1] = val; break; + } +} + +cbuffer cbuffer_uniforms : register(b4, space1) { + uint4 uniforms[1]; +}; +static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); + +[numthreads(1, 1, 1)] +void main() { + set_vector_float2x4(m1, uniforms[0].x, float4((1.0f).xxxx)); + return; +} diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.msl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.msl new file mode 100644 index 0000000000..6c3bb03485 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + thread float2x4 tint_symbol_1 = float2x4(0.0f); + tint_symbol_1[(*(tint_symbol_2)).i] = float4(1.0f); + return; +} + diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.spvasm b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.spvasm new file mode 100644 index 0000000000..f66926472b --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.spvasm @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %m1 "m1" + OpName %main "main" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 +%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float + %10 = OpConstantNull %mat2v4float + %m1 = OpVariable %_ptr_Private_mat2v4float Private %10 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Private_v4float = OpTypePointer Private %v4float + %float_1 = OpConstant %float 1 + %22 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1 + %main = OpFunction %void None %11 + %14 = OpLabel + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %20 = OpAccessChain %_ptr_Private_v4float %m1 %18 + OpStore %20 %22 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.wgsl b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.wgsl new file mode 100644 index 0000000000..acafd9d967 --- /dev/null +++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +[[block]] +struct Uniforms { + i : u32; + j : u32; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +var m1 : mat2x4; + +[[stage(compute), workgroup_size(1)]] +fn main() { + m1[uniforms.i] = vec4(1.0); +} diff --git a/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.spvasm.expected.hlsl deleted file mode 100644 index f78f210ae7..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,218 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_25 : register(b0, space0) { - uint4 x_25[1]; -}; -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -float3 drawShape_vf2_(inout float2 pos) { - bool c2 = false; - bool c3 = false; - bool c4 = false; - bool c5 = false; - bool c6 = false; - int GLF_live4i = 0; - int GLF_live4_looplimiter5 = 0; - float4x2 GLF_live7m42 = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x3 GLF_live7m33 = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int GLF_live7cols = 0; - int GLF_live7_looplimiter3 = 0; - int GLF_live7rows = 0; - int GLF_live7_looplimiter2 = 0; - int GLF_live7_looplimiter1 = 0; - int GLF_live7c = 0; - int GLF_live7r = 0; - int GLF_live7_looplimiter0 = 0; - int GLF_live7sum_index = 0; - int GLF_live7_looplimiter7 = 0; - int GLF_live7cols_1 = 0; - int GLF_live7rows_1 = 0; - float GLF_live7sums[9] = (float[9])0; - int GLF_live7c_1 = 0; - int GLF_live7r_1 = 0; - int x_180 = 0; - float3x3 indexable = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const float x_182 = pos.x; - c2 = (x_182 > 1.0f); - if (c2) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_188 = pos.y; - c3 = (x_188 < 1.0f); - if (c3) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_194 = pos.y; - c4 = (x_194 > 1.0f); - if (c4) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_200 = pos.x; - c5 = (x_200 < 1.0f); - if (c5) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_206 = pos.x; - c6 = ((x_206 + 1.0f) > 1.0f); - if (c6) { - return float3(1.0f, 1.0f, 1.0f); - } - GLF_live4i = 0; - { - [loop] for(; (GLF_live4i < 4); GLF_live4i = (GLF_live4i + 1)) { - if ((GLF_live4_looplimiter5 >= 7)) { - break; - } - GLF_live4_looplimiter5 = (GLF_live4_looplimiter5 + 1); - GLF_live7m42 = float4x2(float2(1.0f, 0.0f), float2(0.0f, 1.0f), float2(0.0f, 0.0f), float2(1.0f, 0.0f)); - GLF_live7m33 = float3x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f)); - GLF_live7cols = 2; - { - [loop] for(; (GLF_live7cols < 4); GLF_live7cols = (GLF_live7cols + 1)) { - if ((GLF_live7_looplimiter3 >= 7)) { - break; - } - GLF_live7_looplimiter3 = (GLF_live7_looplimiter3 + 1); - GLF_live7rows = 2; - { - [loop] for(; (GLF_live7rows < 4); GLF_live7rows = (GLF_live7rows + 1)) { - if ((GLF_live7_looplimiter2 >= 7)) { - break; - } - GLF_live7_looplimiter2 = (GLF_live7_looplimiter2 + 1); - GLF_live7_looplimiter1 = 0; - GLF_live7c = 0; - { - [loop] for(; (GLF_live7c < 3); GLF_live7c = (GLF_live7c + 1)) { - if ((GLF_live7_looplimiter1 >= 7)) { - break; - } - GLF_live7_looplimiter1 = (GLF_live7_looplimiter1 + 1); - GLF_live7r = 0; - { - [loop] for(; (GLF_live7r < 2); GLF_live7r = (GLF_live7r + 1)) { - if ((GLF_live7_looplimiter0 >= 7)) { - break; - } - GLF_live7_looplimiter0 = (GLF_live7_looplimiter0 + 1); - set_float3(GLF_live7m33[(((GLF_live7c >= 0) & (GLF_live7c < 3)) ? GLF_live7c : 0)], (((GLF_live7r >= 0) & (GLF_live7r < 3)) ? GLF_live7r : 0), 1.0f); - const float x_267 = asfloat(x_25[0].y); - if ((0.0f > x_267)) { - } else { - set_float2(GLF_live7m42[(((GLF_live7c >= 0) & (GLF_live7c < 4)) ? GLF_live7c : 0)], (((GLF_live7r >= 0) & (GLF_live7r < 2)) ? GLF_live7r : 0), 1.0f); - } - } - } - } - } - } - } - } - } - GLF_live7sum_index = 0; - GLF_live7_looplimiter7 = 0; - GLF_live7cols_1 = 2; - { - [loop] for(; (GLF_live7cols_1 < 4); GLF_live7cols_1 = (GLF_live7cols_1 + 1)) { - if ((GLF_live7_looplimiter7 >= 7)) { - break; - } - GLF_live7_looplimiter7 = (GLF_live7_looplimiter7 + 1); - GLF_live7rows_1 = 2; - GLF_live7sums[(((GLF_live7sum_index >= 0) & (GLF_live7sum_index < 9)) ? GLF_live7sum_index : 0)] = 0.0f; - GLF_live7c_1 = 0; - { - [loop] for(; (GLF_live7c_1 < 1); GLF_live7c_1 = (GLF_live7c_1 + 1)) { - GLF_live7r_1 = 0; - { - [loop] for(; (GLF_live7r_1 < GLF_live7rows_1); GLF_live7r_1 = (GLF_live7r_1 + 1)) { - const int x_310 = (((GLF_live7sum_index >= 0) & (GLF_live7sum_index < 9)) ? GLF_live7sum_index : 0); - const float3x3 x_312 = transpose(GLF_live7m33); - if ((GLF_live7c_1 < 3)) { - x_180 = 1; - } else { - const float x_318 = asfloat(x_25[0].x); - x_180 = int(x_318); - } - const int x_320 = x_180; - const int x_93 = GLF_live7r_1; - indexable = x_312; - const float x_324 = indexable[x_320][((x_93 < 3) ? 1 : 0)]; - const float x_326 = GLF_live7sums[x_310]; - GLF_live7sums[x_310] = (x_326 + x_324); - const int x_332 = (((GLF_live7sum_index >= 0) & (GLF_live7sum_index < 9)) ? GLF_live7sum_index : 0); - const float x_334 = GLF_live7m42[1][GLF_live7r_1]; - const float x_336 = GLF_live7sums[x_332]; - GLF_live7sums[x_332] = (x_336 + x_334); - } - } - } - } - GLF_live7sum_index = (GLF_live7sum_index + 1); - } - } - } - } - return float3(1.0f, 1.0f, 1.0f); -} - -void main_1() { - float2 position = float2(0.0f, 0.0f); - float2 param = float2(0.0f, 0.0f); - float2 param_1 = float2(0.0f, 0.0f); - int i = 0; - float2 param_2 = float2(0.0f, 0.0f); - const float x_161 = asfloat(x_25[0].x); - if ((x_161 >= 2.0f)) { - const float4 x_165 = gl_FragCoord; - position = float2(x_165.x, x_165.y); - param = position; - const float3 x_168 = drawShape_vf2_(param); - param_1 = position; - const float3 x_170 = drawShape_vf2_(param_1); - i = 25; - { - [loop] for(; (i > 0); i = (i - 1)) { - param_2 = position; - const float3 x_178 = drawShape_vf2_(param_2); - } - } - } - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000002DBA887C060(100,28-81): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 6b45ab873b..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/control-flow-in-function/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,246 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_25 : register(b0, space0) { - uint4 x_25[1]; -}; -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -float3 drawShape_vf2_(inout float2 pos) { - bool c2 = false; - bool c3 = false; - bool c4 = false; - bool c5 = false; - bool c6 = false; - int GLF_live4i = 0; - int GLF_live4_looplimiter5 = 0; - float4x2 GLF_live7m42 = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x3 GLF_live7m33 = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int GLF_live7cols = 0; - int GLF_live7_looplimiter3 = 0; - int GLF_live7rows = 0; - int GLF_live7_looplimiter2 = 0; - int GLF_live7_looplimiter1 = 0; - int GLF_live7c = 0; - int GLF_live7r = 0; - int GLF_live7_looplimiter0 = 0; - int GLF_live7sum_index = 0; - int GLF_live7_looplimiter7 = 0; - int GLF_live7cols_1 = 0; - int GLF_live7rows_1 = 0; - float GLF_live7sums[9] = (float[9])0; - int GLF_live7c_1 = 0; - int GLF_live7r_1 = 0; - int x_180 = 0; - float3x3 indexable = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const float x_182 = pos.x; - c2 = (x_182 > 1.0f); - if (c2) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_188 = pos.y; - c3 = (x_188 < 1.0f); - if (c3) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_194 = pos.y; - c4 = (x_194 > 1.0f); - if (c4) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_200 = pos.x; - c5 = (x_200 < 1.0f); - if (c5) { - return float3(1.0f, 1.0f, 1.0f); - } - const float x_206 = pos.x; - c6 = ((x_206 + 1.0f) > 1.0f); - if (c6) { - return float3(1.0f, 1.0f, 1.0f); - } - GLF_live4i = 0; - { - [loop] for(; (GLF_live4i < 4); GLF_live4i = (GLF_live4i + 1)) { - if ((GLF_live4_looplimiter5 >= 7)) { - break; - } - GLF_live4_looplimiter5 = (GLF_live4_looplimiter5 + 1); - GLF_live7m42 = float4x2(float2(1.0f, 0.0f), float2(0.0f, 1.0f), float2(0.0f, 0.0f), float2(1.0f, 0.0f)); - GLF_live7m33 = float3x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f)); - GLF_live7cols = 2; - { - [loop] for(; (GLF_live7cols < 4); GLF_live7cols = (GLF_live7cols + 1)) { - if ((GLF_live7_looplimiter3 >= 7)) { - break; - } - GLF_live7_looplimiter3 = (GLF_live7_looplimiter3 + 1); - GLF_live7rows = 2; - { - [loop] for(; (GLF_live7rows < 4); GLF_live7rows = (GLF_live7rows + 1)) { - if ((GLF_live7_looplimiter2 >= 7)) { - break; - } - GLF_live7_looplimiter2 = (GLF_live7_looplimiter2 + 1); - GLF_live7_looplimiter1 = 0; - GLF_live7c = 0; - { - [loop] for(; (GLF_live7c < 3); GLF_live7c = (GLF_live7c + 1)) { - if ((GLF_live7_looplimiter1 >= 7)) { - break; - } - GLF_live7_looplimiter1 = (GLF_live7_looplimiter1 + 1); - GLF_live7r = 0; - { - [loop] for(; (GLF_live7r < 2); GLF_live7r = (GLF_live7r + 1)) { - if ((GLF_live7_looplimiter0 >= 7)) { - break; - } - GLF_live7_looplimiter0 = (GLF_live7_looplimiter0 + 1); - bool tint_tmp = (GLF_live7c >= 0); - if (tint_tmp) { - tint_tmp = (GLF_live7c < 3); - } - bool tint_tmp_1 = (GLF_live7r >= 0); - if (tint_tmp_1) { - tint_tmp_1 = (GLF_live7r < 3); - } - set_float3(GLF_live7m33[((tint_tmp) ? GLF_live7c : 0)], ((tint_tmp_1) ? GLF_live7r : 0), 1.0f); - const float x_267 = asfloat(x_25[0].y); - if ((0.0f > x_267)) { - } else { - bool tint_tmp_2 = (GLF_live7c >= 0); - if (tint_tmp_2) { - tint_tmp_2 = (GLF_live7c < 4); - } - bool tint_tmp_3 = (GLF_live7r >= 0); - if (tint_tmp_3) { - tint_tmp_3 = (GLF_live7r < 2); - } - set_float2(GLF_live7m42[((tint_tmp_2) ? GLF_live7c : 0)], ((tint_tmp_3) ? GLF_live7r : 0), 1.0f); - } - } - } - } - } - } - } - } - } - GLF_live7sum_index = 0; - GLF_live7_looplimiter7 = 0; - GLF_live7cols_1 = 2; - { - [loop] for(; (GLF_live7cols_1 < 4); GLF_live7cols_1 = (GLF_live7cols_1 + 1)) { - if ((GLF_live7_looplimiter7 >= 7)) { - break; - } - GLF_live7_looplimiter7 = (GLF_live7_looplimiter7 + 1); - GLF_live7rows_1 = 2; - bool tint_tmp_4 = (GLF_live7sum_index >= 0); - if (tint_tmp_4) { - tint_tmp_4 = (GLF_live7sum_index < 9); - } - GLF_live7sums[((tint_tmp_4) ? GLF_live7sum_index : 0)] = 0.0f; - GLF_live7c_1 = 0; - { - [loop] for(; (GLF_live7c_1 < 1); GLF_live7c_1 = (GLF_live7c_1 + 1)) { - GLF_live7r_1 = 0; - { - [loop] for(; (GLF_live7r_1 < GLF_live7rows_1); GLF_live7r_1 = (GLF_live7r_1 + 1)) { - bool tint_tmp_5 = (GLF_live7sum_index >= 0); - if (tint_tmp_5) { - tint_tmp_5 = (GLF_live7sum_index < 9); - } - const int x_310 = ((tint_tmp_5) ? GLF_live7sum_index : 0); - const float3x3 x_312 = transpose(GLF_live7m33); - if ((GLF_live7c_1 < 3)) { - x_180 = 1; - } else { - const float x_318 = asfloat(x_25[0].x); - x_180 = int(x_318); - } - const int x_320 = x_180; - const int x_93 = GLF_live7r_1; - indexable = x_312; - const float x_324 = indexable[x_320][((x_93 < 3) ? 1 : 0)]; - const float x_326 = GLF_live7sums[x_310]; - GLF_live7sums[x_310] = (x_326 + x_324); - bool tint_tmp_6 = (GLF_live7sum_index >= 0); - if (tint_tmp_6) { - tint_tmp_6 = (GLF_live7sum_index < 9); - } - const int x_332 = ((tint_tmp_6) ? GLF_live7sum_index : 0); - const float x_334 = GLF_live7m42[1][GLF_live7r_1]; - const float x_336 = GLF_live7sums[x_332]; - GLF_live7sums[x_332] = (x_336 + x_334); - } - } - } - } - GLF_live7sum_index = (GLF_live7sum_index + 1); - } - } - } - } - return float3(1.0f, 1.0f, 1.0f); -} - -void main_1() { - float2 position = float2(0.0f, 0.0f); - float2 param = float2(0.0f, 0.0f); - float2 param_1 = float2(0.0f, 0.0f); - int i = 0; - float2 param_2 = float2(0.0f, 0.0f); - const float x_161 = asfloat(x_25[0].x); - if ((x_161 >= 2.0f)) { - const float4 x_165 = gl_FragCoord; - position = float2(x_165.x, x_165.y); - param = position; - const float3 x_168 = drawShape_vf2_(param); - param_1 = position; - const float3 x_170 = drawShape_vf2_(param_1); - i = 25; - { - [loop] for(; (i > 0); i = (i - 1)) { - param_2 = position; - const float3 x_178 = drawShape_vf2_(param_2); - } - } - } - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000018D92558420(100,28-81): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.spvasm.expected.hlsl deleted file mode 100644 index 909cf6f032..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,80 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -cbuffer cbuffer_x_6 : register(b0, space0) { - uint4 x_6[4]; -}; -cbuffer cbuffer_x_9 : register(b1, space0) { - uint4 x_9[2]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float3x3 m = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int a = 0; - float3 arr[2] = (float3[2])0; - float3 v = float3(0.0f, 0.0f, 0.0f); - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_45 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - const float x_46 = float(x_45); - m = float3x3(float3(x_46, 0.0f, 0.0f), float3(0.0f, x_46, 0.0f), float3(0.0f, 0.0f, x_46)); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_52 = asint(x_6[scalar_offset_1 / 4][scalar_offset_1 % 4]); - a = x_52; - const int x_53 = a; - const int x_54 = a; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_56 = asfloat(x_9[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float3(m[x_53], x_54, x_56); - const float3 x_59 = m[1]; - const float3 x_61 = m[1]; - const float3 tint_symbol_3[2] = {x_59, x_61}; - arr = tint_symbol_3; - const float x_64 = asfloat(x_9[1].x); - v = float3(x_64, x_64, x_64); - const float3 x_68 = arr[a]; - v = (v + x_68); - const float3 x_71 = v; - const int x_73 = asint(x_6[1].x); - const int x_76 = asint(x_6[2].x); - const int x_79 = asint(x_6[1].x); - if (all((x_71 == float3(float(x_73), float(x_76), float(x_79))))) { - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_88 = asint(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const int x_91 = asint(x_6[3].x); - const int x_94 = asint(x_6[3].x); - const uint scalar_offset_4 = ((16u * uint(0))) / 4; - const int x_97 = asint(x_6[scalar_offset_4 / 4][scalar_offset_4 % 4]); - x_GLF_color = float4(float(x_88), float(x_91), float(x_94), float(x_97)); - } else { - const int x_101 = asint(x_6[3].x); - const float x_102 = float(x_101); - x_GLF_color = float4(x_102, x_102, x_102, x_102); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x00000167B1373790(29,14-20): error X3500: array reference cannot be used as an l-value; not natively addressable - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 83a7b7d3f9..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-array-matrix-element/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,80 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -cbuffer cbuffer_x_6 : register(b0, space0) { - uint4 x_6[4]; -}; -cbuffer cbuffer_x_9 : register(b1, space0) { - uint4 x_9[2]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float3x3 m = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int a = 0; - float3 arr[2] = (float3[2])0; - float3 v = float3(0.0f, 0.0f, 0.0f); - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_45 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - const float x_46 = float(x_45); - m = float3x3(float3(x_46, 0.0f, 0.0f), float3(0.0f, x_46, 0.0f), float3(0.0f, 0.0f, x_46)); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_52 = asint(x_6[scalar_offset_1 / 4][scalar_offset_1 % 4]); - a = x_52; - const int x_53 = a; - const int x_54 = a; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_56 = asfloat(x_9[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float3(m[x_53], x_54, x_56); - const float3 x_59 = m[1]; - const float3 x_61 = m[1]; - const float3 tint_symbol_3[2] = {x_59, x_61}; - arr = tint_symbol_3; - const float x_64 = asfloat(x_9[1].x); - v = float3(x_64, x_64, x_64); - const float3 x_68 = arr[a]; - v = (v + x_68); - const float3 x_71 = v; - const int x_73 = asint(x_6[1].x); - const int x_76 = asint(x_6[2].x); - const int x_79 = asint(x_6[1].x); - if (all((x_71 == float3(float(x_73), float(x_76), float(x_79))))) { - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_88 = asint(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const int x_91 = asint(x_6[3].x); - const int x_94 = asint(x_6[3].x); - const uint scalar_offset_4 = ((16u * uint(0))) / 4; - const int x_97 = asint(x_6[scalar_offset_4 / 4][scalar_offset_4 % 4]); - x_GLF_color = float4(float(x_88), float(x_91), float(x_94), float(x_97)); - } else { - const int x_101 = asint(x_6[3].x); - const float x_102 = float(x_101); - x_GLF_color = float4(x_102, x_102, x_102, x_102); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001E84DF5F800(29,14-20): error X3500: array reference cannot be used as an l-value; not natively addressable - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.spvasm.expected.hlsl deleted file mode 100644 index e2a6cb32e3..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,78 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -cbuffer cbuffer_x_6 : register(b0, space0) { - uint4 x_6[2]; -}; -cbuffer cbuffer_x_9 : register(b1, space0) { - uint4 x_9[4]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - int a = 0; - float4 v = float4(0.0f, 0.0f, 0.0f, 0.0f); - float3x4 m = float3x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x4 indexable = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_44 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - a = x_44; - const float x_46 = asfloat(x_9[2].x); - v = float4(x_46, x_46, x_46, x_46); - const float x_49 = asfloat(x_9[3].x); - m = float3x4(float4(x_49, 0.0f, 0.0f, 0.0f), float4(0.0f, x_49, 0.0f, 0.0f), float4(0.0f, 0.0f, x_49, 0.0f)); - const int x_54 = a; - const int x_55 = a; - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_57 = asfloat(x_9[scalar_offset_1 / 4][scalar_offset_1 % 4]); - set_float4(m[x_54], x_55, x_57); - const int x_59 = a; - const float3x4 x_60 = m; - const int x_78 = a; - const int x_79 = a; - indexable = float4x4(float4(x_60[0u].x, x_60[0u].y, x_60[0u].z, x_60[0u].w), float4(x_60[1u].x, x_60[1u].y, x_60[1u].z, x_60[1u].w), float4(x_60[2u].x, x_60[2u].y, x_60[2u].z, x_60[2u].w), float4(0.0f, 0.0f, 0.0f, 1.0f)); - const float x_81 = indexable[x_78][x_79]; - const float x_83 = v[x_59]; - set_float4(v, x_59, (x_83 + x_81)); - const float x_87 = v.y; - const float x_89 = asfloat(x_9[1].x); - if ((x_87 == x_89)) { - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const int x_95 = asint(x_6[scalar_offset_2 / 4][scalar_offset_2 % 4]); - const int x_98 = asint(x_6[1].x); - const int x_101 = asint(x_6[1].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_104 = asint(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); - x_GLF_color = float4(float(x_95), float(x_98), float(x_101), float(x_104)); - } else { - const int x_108 = asint(x_6[1].x); - const float x_109 = float(x_108); - x_GLF_color = float4(x_109, x_109, x_109, x_109); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_3 = {x_GLF_color}; - return tint_symbol_3; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000015D0949C000(29,14-20): error X3500: array reference cannot be used as an l-value; not natively addressable - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 884c3d6b83..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-increment-vector-component-with-matrix-copy/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,78 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -cbuffer cbuffer_x_6 : register(b0, space0) { - uint4 x_6[2]; -}; -cbuffer cbuffer_x_9 : register(b1, space0) { - uint4 x_9[4]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - int a = 0; - float4 v = float4(0.0f, 0.0f, 0.0f, 0.0f); - float3x4 m = float3x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x4 indexable = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_44 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - a = x_44; - const float x_46 = asfloat(x_9[2].x); - v = float4(x_46, x_46, x_46, x_46); - const float x_49 = asfloat(x_9[3].x); - m = float3x4(float4(x_49, 0.0f, 0.0f, 0.0f), float4(0.0f, x_49, 0.0f, 0.0f), float4(0.0f, 0.0f, x_49, 0.0f)); - const int x_54 = a; - const int x_55 = a; - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_57 = asfloat(x_9[scalar_offset_1 / 4][scalar_offset_1 % 4]); - set_float4(m[x_54], x_55, x_57); - const int x_59 = a; - const float3x4 x_60 = m; - const int x_78 = a; - const int x_79 = a; - indexable = float4x4(float4(x_60[0u].x, x_60[0u].y, x_60[0u].z, x_60[0u].w), float4(x_60[1u].x, x_60[1u].y, x_60[1u].z, x_60[1u].w), float4(x_60[2u].x, x_60[2u].y, x_60[2u].z, x_60[2u].w), float4(0.0f, 0.0f, 0.0f, 1.0f)); - const float x_81 = indexable[x_78][x_79]; - const float x_83 = v[x_59]; - set_float4(v, x_59, (x_83 + x_81)); - const float x_87 = v.y; - const float x_89 = asfloat(x_9[1].x); - if ((x_87 == x_89)) { - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const int x_95 = asint(x_6[scalar_offset_2 / 4][scalar_offset_2 % 4]); - const int x_98 = asint(x_6[1].x); - const int x_101 = asint(x_6[1].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_104 = asint(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); - x_GLF_color = float4(float(x_95), float(x_98), float(x_101), float(x_104)); - } else { - const int x_108 = asint(x_6[1].x); - const float x_109 = float(x_108); - x_GLF_color = float4(x_109, x_109, x_109, x_109); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_3 = {x_GLF_color}; - return tint_symbol_3; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000002C75C7A2100(29,14-20): error X3500: array reference cannot be used as an l-value; not natively addressable - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.spvasm.expected.hlsl deleted file mode 100644 index 097dd84ffb..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,105 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -cbuffer cbuffer_x_7 : register(b1, space0) { - uint4 x_7[2]; -}; -cbuffer cbuffer_x_10 : register(b0, space0) { - uint4 x_10[4]; -}; -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x3 m23 = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int i = 0; - const float x_46 = asfloat(x_7[1].x); - m23 = float2x3(float3(x_46, 0.0f, 0.0f), float3(0.0f, x_46, 0.0f)); - i = 1; - [loop] while (true) { - bool x_80 = false; - bool x_81_phi = false; - const int x_54 = i; - const int x_56 = asint(x_10[3].x); - if ((x_54 < x_56)) { - } else { - break; - } - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_60 = asint(x_10[scalar_offset / 4][scalar_offset % 4]); - const int x_62 = asint(x_10[2].x); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_64 = asfloat(x_7[scalar_offset_1 / 4][scalar_offset_1 % 4]); - const float x_66 = m23[x_60][x_62]; - set_float3(m23[x_60], x_62, (x_66 + x_64)); - const float x_70 = gl_FragCoord.y; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_72 = asfloat(x_7[scalar_offset_2 / 4][scalar_offset_2 % 4]); - if ((x_70 < x_72)) { - } - x_81_phi = true; - if (true) { - const float x_79 = gl_FragCoord.x; - x_80 = (x_79 < 0.0f); - x_81_phi = x_80; - } - if (!(x_81_phi)) { - break; - } - { - i = (i + 1); - } - } - const float2x3 x_87 = m23; - const int x_89 = asint(x_10[1].x); - const int x_92 = asint(x_10[1].x); - const int x_95 = asint(x_10[1].x); - const int x_98 = asint(x_10[1].x); - const int x_101 = asint(x_10[1].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_104 = asint(x_10[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const float2x3 x_108 = float2x3(float3(float(x_89), float(x_92), float(x_95)), float3(float(x_98), float(x_101), float(x_104))); - if ((all((x_87[0u] == x_108[0u])) & all((x_87[1u] == x_108[1u])))) { - const uint scalar_offset_4 = ((16u * uint(0))) / 4; - const int x_122 = asint(x_10[scalar_offset_4 / 4][scalar_offset_4 % 4]); - const int x_125 = asint(x_10[1].x); - const int x_128 = asint(x_10[1].x); - const uint scalar_offset_5 = ((16u * uint(0))) / 4; - const int x_131 = asint(x_10[scalar_offset_5 / 4][scalar_offset_5 % 4]); - x_GLF_color = float4(float(x_122), float(x_125), float(x_128), float(x_131)); - } else { - const int x_135 = asint(x_10[1].x); - const float x_136 = float(x_135); - x_GLF_color = float4(x_136, x_136, x_136, x_136); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_5 = {x_GLF_color}; - return tint_symbol_5; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001FC5E0CC510(20,10-21): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.wgsl.expected.hlsl deleted file mode 100644 index eb7806c2f1..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-matrix-element-break-after-first-iteration/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,109 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -cbuffer cbuffer_x_7 : register(b1, space0) { - uint4 x_7[2]; -}; -cbuffer cbuffer_x_10 : register(b0, space0) { - uint4 x_10[4]; -}; -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x3 m23 = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int i = 0; - const float x_46 = asfloat(x_7[1].x); - m23 = float2x3(float3(x_46, 0.0f, 0.0f), float3(0.0f, x_46, 0.0f)); - i = 1; - [loop] while (true) { - bool x_80 = false; - bool x_81_phi = false; - const int x_54 = i; - const int x_56 = asint(x_10[3].x); - if ((x_54 < x_56)) { - } else { - break; - } - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_60 = asint(x_10[scalar_offset / 4][scalar_offset % 4]); - const int x_62 = asint(x_10[2].x); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_64 = asfloat(x_7[scalar_offset_1 / 4][scalar_offset_1 % 4]); - const float x_66 = m23[x_60][x_62]; - set_float3(m23[x_60], x_62, (x_66 + x_64)); - const float x_70 = gl_FragCoord.y; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_72 = asfloat(x_7[scalar_offset_2 / 4][scalar_offset_2 % 4]); - if ((x_70 < x_72)) { - } - x_81_phi = true; - if (true) { - const float x_79 = gl_FragCoord.x; - x_80 = (x_79 < 0.0f); - x_81_phi = x_80; - } - if (!(x_81_phi)) { - break; - } - { - i = (i + 1); - } - } - const float2x3 x_87 = m23; - const int x_89 = asint(x_10[1].x); - const int x_92 = asint(x_10[1].x); - const int x_95 = asint(x_10[1].x); - const int x_98 = asint(x_10[1].x); - const int x_101 = asint(x_10[1].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_104 = asint(x_10[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const float2x3 x_108 = float2x3(float3(float(x_89), float(x_92), float(x_95)), float3(float(x_98), float(x_101), float(x_104))); - bool tint_tmp = all((x_87[0u] == x_108[0u])); - if (tint_tmp) { - tint_tmp = all((x_87[1u] == x_108[1u])); - } - if ((tint_tmp)) { - const uint scalar_offset_4 = ((16u * uint(0))) / 4; - const int x_122 = asint(x_10[scalar_offset_4 / 4][scalar_offset_4 % 4]); - const int x_125 = asint(x_10[1].x); - const int x_128 = asint(x_10[1].x); - const uint scalar_offset_5 = ((16u * uint(0))) / 4; - const int x_131 = asint(x_10[scalar_offset_5 / 4][scalar_offset_5 % 4]); - x_GLF_color = float4(float(x_122), float(x_125), float(x_128), float(x_131)); - } else { - const int x_135 = asint(x_10[1].x); - const float x_136 = float(x_135); - x_GLF_color = float4(x_136, x_136, x_136, x_136); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_5 = {x_GLF_color}; - return tint_symbol_5; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000017C45B3F9B0(20,10-21): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.spvasm.expected.hlsl deleted file mode 100644 index cfed15c803..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,334 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -static int x_GLF_global_loop_count = 0; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x3 m23 = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float2x4 m24 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x2 m32 = float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x3 m33 = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x4 m34 = float3x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x2 m42 = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x3 m43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x4 m44 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int i = 0; - int i_1 = 0; - int i_2 = 0; - int i_3 = 0; - int i_4 = 0; - int i_5 = 0; - int i_6 = 0; - int i_7 = 0; - int i_8 = 0; - int i_9 = 0; - int i_10 = 0; - int i_11 = 0; - int i_12 = 0; - int i_13 = 0; - int i_14 = 0; - int i_15 = 0; - int i_16 = 0; - int i_17 = 0; - int i_18 = 0; - int i_19 = 0; - int i_20 = 0; - int i_21 = 0; - int i_22 = 0; - int i_23 = 0; - int i_24 = 0; - int i_25 = 0; - int i_26 = 0; - int i_27 = 0; - int i_28 = 0; - int i_29 = 0; - int i_30 = 0; - int i_31 = 0; - int i_32 = 0; - int i_33 = 0; - int i_34 = 0; - int i_35 = 0; - int i_36 = 0; - int i_37 = 0; - float sum = 0.0f; - int r = 0; - x_GLF_global_loop_count = 0; - m23 = float2x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m24 = float2x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - m32 = float3x2(float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f)); - m33 = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m34 = float3x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - m42 = float4x2(float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f)); - m43 = float4x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m44 = float4x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - i = 0; - { - [loop] for(; (i < 1); i = (i + 1)) { - i_1 = 0; - { - [loop] for(; (i_1 < 1); i_1 = (i_1 + 1)) { - i_2 = 0; - { - [loop] for(; (i_2 < 1); i_2 = (i_2 + 1)) { - i_3 = 0; - { - [loop] for(; (i_3 < 1); i_3 = (i_3 + 1)) { - i_4 = 0; - { - [loop] for(; (i_4 < 1); i_4 = (i_4 + 1)) { - i_5 = 0; - { - [loop] for(; (i_5 < 1); i_5 = (i_5 + 1)) { - i_6 = 0; - { - [loop] for(; (i_6 < 1); i_6 = (i_6 + 1)) { - i_7 = 0; - { - [loop] for(; (i_7 < 1); i_7 = (i_7 + 1)) { - i_8 = 0; - { - [loop] for(; (i_8 < 1); i_8 = (i_8 + 1)) { - i_9 = 0; - { - [loop] for(; (i_9 < 1); i_9 = (i_9 + 1)) { - i_10 = 0; - { - [loop] for(; (i_10 < 1); i_10 = (i_10 + 1)) { - i_11 = 0; - { - [loop] for(; (i_11 < 1); i_11 = (i_11 + 1)) { - i_12 = 0; - { - [loop] for(; (i_12 < 1); i_12 = (i_12 + 1)) { - i_13 = 0; - { - [loop] for(; (i_13 < 1); i_13 = (i_13 + 1)) { - i_14 = 0; - { - [loop] for(; (i_14 < 1); i_14 = (i_14 + 1)) { - i_15 = 0; - { - [loop] for(; (i_15 < 1); i_15 = (i_15 + 1)) { - i_16 = 0; - { - [loop] for(; (i_16 < 1); i_16 = (i_16 + 1)) { - i_17 = 0; - { - [loop] for(; (i_17 < 1); i_17 = (i_17 + 1)) { - i_18 = 0; - { - [loop] for(; (i_18 < 1); i_18 = (i_18 + 1)) { - i_19 = 0; - { - [loop] for(; (i_19 < 1); i_19 = (i_19 + 1)) { - i_20 = 0; - { - [loop] for(; (i_20 < 1); i_20 = (i_20 + 1)) { - i_21 = 0; - { - [loop] for(; (i_21 < 1); i_21 = (i_21 + 1)) { - i_22 = 0; - { - [loop] for(; (i_22 < 1); i_22 = (i_22 + 1)) { - i_23 = 0; - { - [loop] for(; (i_23 < 1); i_23 = (i_23 + 1)) { - i_24 = 0; - { - [loop] for(; (i_24 < 1); i_24 = (i_24 + 1)) { - i_25 = 0; - { - [loop] for(; (i_25 < 1); i_25 = (i_25 + 1)) { - i_26 = 0; - { - [loop] for(; (i_26 < 1); i_26 = (i_26 + 1)) { - i_27 = 0; - { - [loop] for(; (i_27 < 1); i_27 = (i_27 + 1)) { - i_28 = 0; - { - [loop] for(; (i_28 < 1); i_28 = (i_28 + 1)) { - i_29 = 0; - { - [loop] for(; (i_29 < 1); i_29 = (i_29 + 1)) { - i_30 = 0; - { - [loop] for(; (i_30 < 1); i_30 = (i_30 + 1)) { - i_31 = 0; - { - [loop] for(; (i_31 < 1); i_31 = (i_31 + 1)) { - i_32 = 0; - { - [loop] for(; (i_32 < 1); i_32 = (i_32 + 1)) { - i_33 = 0; - { - [loop] for(; (i_33 < 1); i_33 = (i_33 + 1)) { - i_34 = 0; - { - [loop] for(; (i_34 < 1); i_34 = (i_34 + 1)) { - i_35 = 0; - { - [loop] for(; (i_35 < 1); i_35 = (i_35 + 1)) { - i_36 = 0; - { - [loop] for(; (i_36 < 1); i_36 = (i_36 + 1)) { - i_37 = 0; - { - [loop] for(; (i_37 < 1); i_37 = (i_37 + 1)) { - [loop] while (true) { - x_GLF_global_loop_count = (x_GLF_global_loop_count + 1); - { - if ((x_GLF_global_loop_count < 98)) { - } else { - break; - } - } - } - set_float3(m23[i_37], i_37, 1.0f); - set_float4(m24[i_37], i_37, 1.0f); - set_float2(m32[i_37], i_37, 1.0f); - set_float3(m33[i_37], i_37, 1.0f); - set_float4(m34[i_37], i_37, 1.0f); - set_float2(m42[i_37], i_37, 1.0f); - set_float3(m43[i_37], i_37, 1.0f); - set_float4(m44[i_37], i_37, 1.0f); - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - sum = 0.0f; - r = 0; - { - [loop] for(; (x_GLF_global_loop_count < 100); r = (r + 1)) { - x_GLF_global_loop_count = (x_GLF_global_loop_count + 1); - const float x_486 = m23[0][r]; - sum = (sum + x_486); - const float x_491 = m24[0][r]; - sum = (sum + x_491); - const float x_496 = m32[0][r]; - sum = (sum + x_496); - const float x_501 = m33[0][r]; - sum = (sum + x_501); - const float x_506 = m34[0][r]; - sum = (sum + x_506); - const float x_511 = m42[0][r]; - sum = (sum + x_511); - const float x_516 = m43[0][r]; - sum = (sum + x_516); - const float x_521 = m44[0][r]; - sum = (sum + x_521); - } - } - if ((sum == 8.0f)) { - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - } else { - x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_1 = {x_GLF_color}; - return tint_symbol_1; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001DD183F6C40(187,160-195): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 1b87903284..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-multiple-one-iteration-loops-global-counter-write-matrices/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,334 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -static int x_GLF_global_loop_count = 0; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x3 m23 = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float2x4 m24 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x2 m32 = float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x3 m33 = float3x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float3x4 m34 = float3x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x2 m42 = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x3 m43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float4x4 m44 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int i = 0; - int i_1 = 0; - int i_2 = 0; - int i_3 = 0; - int i_4 = 0; - int i_5 = 0; - int i_6 = 0; - int i_7 = 0; - int i_8 = 0; - int i_9 = 0; - int i_10 = 0; - int i_11 = 0; - int i_12 = 0; - int i_13 = 0; - int i_14 = 0; - int i_15 = 0; - int i_16 = 0; - int i_17 = 0; - int i_18 = 0; - int i_19 = 0; - int i_20 = 0; - int i_21 = 0; - int i_22 = 0; - int i_23 = 0; - int i_24 = 0; - int i_25 = 0; - int i_26 = 0; - int i_27 = 0; - int i_28 = 0; - int i_29 = 0; - int i_30 = 0; - int i_31 = 0; - int i_32 = 0; - int i_33 = 0; - int i_34 = 0; - int i_35 = 0; - int i_36 = 0; - int i_37 = 0; - float sum = 0.0f; - int r = 0; - x_GLF_global_loop_count = 0; - m23 = float2x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m24 = float2x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - m32 = float3x2(float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f)); - m33 = float3x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m34 = float3x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - m42 = float4x2(float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f)); - m43 = float4x3(float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f)); - m44 = float4x4(float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)); - i = 0; - { - [loop] for(; (i < 1); i = (i + 1)) { - i_1 = 0; - { - [loop] for(; (i_1 < 1); i_1 = (i_1 + 1)) { - i_2 = 0; - { - [loop] for(; (i_2 < 1); i_2 = (i_2 + 1)) { - i_3 = 0; - { - [loop] for(; (i_3 < 1); i_3 = (i_3 + 1)) { - i_4 = 0; - { - [loop] for(; (i_4 < 1); i_4 = (i_4 + 1)) { - i_5 = 0; - { - [loop] for(; (i_5 < 1); i_5 = (i_5 + 1)) { - i_6 = 0; - { - [loop] for(; (i_6 < 1); i_6 = (i_6 + 1)) { - i_7 = 0; - { - [loop] for(; (i_7 < 1); i_7 = (i_7 + 1)) { - i_8 = 0; - { - [loop] for(; (i_8 < 1); i_8 = (i_8 + 1)) { - i_9 = 0; - { - [loop] for(; (i_9 < 1); i_9 = (i_9 + 1)) { - i_10 = 0; - { - [loop] for(; (i_10 < 1); i_10 = (i_10 + 1)) { - i_11 = 0; - { - [loop] for(; (i_11 < 1); i_11 = (i_11 + 1)) { - i_12 = 0; - { - [loop] for(; (i_12 < 1); i_12 = (i_12 + 1)) { - i_13 = 0; - { - [loop] for(; (i_13 < 1); i_13 = (i_13 + 1)) { - i_14 = 0; - { - [loop] for(; (i_14 < 1); i_14 = (i_14 + 1)) { - i_15 = 0; - { - [loop] for(; (i_15 < 1); i_15 = (i_15 + 1)) { - i_16 = 0; - { - [loop] for(; (i_16 < 1); i_16 = (i_16 + 1)) { - i_17 = 0; - { - [loop] for(; (i_17 < 1); i_17 = (i_17 + 1)) { - i_18 = 0; - { - [loop] for(; (i_18 < 1); i_18 = (i_18 + 1)) { - i_19 = 0; - { - [loop] for(; (i_19 < 1); i_19 = (i_19 + 1)) { - i_20 = 0; - { - [loop] for(; (i_20 < 1); i_20 = (i_20 + 1)) { - i_21 = 0; - { - [loop] for(; (i_21 < 1); i_21 = (i_21 + 1)) { - i_22 = 0; - { - [loop] for(; (i_22 < 1); i_22 = (i_22 + 1)) { - i_23 = 0; - { - [loop] for(; (i_23 < 1); i_23 = (i_23 + 1)) { - i_24 = 0; - { - [loop] for(; (i_24 < 1); i_24 = (i_24 + 1)) { - i_25 = 0; - { - [loop] for(; (i_25 < 1); i_25 = (i_25 + 1)) { - i_26 = 0; - { - [loop] for(; (i_26 < 1); i_26 = (i_26 + 1)) { - i_27 = 0; - { - [loop] for(; (i_27 < 1); i_27 = (i_27 + 1)) { - i_28 = 0; - { - [loop] for(; (i_28 < 1); i_28 = (i_28 + 1)) { - i_29 = 0; - { - [loop] for(; (i_29 < 1); i_29 = (i_29 + 1)) { - i_30 = 0; - { - [loop] for(; (i_30 < 1); i_30 = (i_30 + 1)) { - i_31 = 0; - { - [loop] for(; (i_31 < 1); i_31 = (i_31 + 1)) { - i_32 = 0; - { - [loop] for(; (i_32 < 1); i_32 = (i_32 + 1)) { - i_33 = 0; - { - [loop] for(; (i_33 < 1); i_33 = (i_33 + 1)) { - i_34 = 0; - { - [loop] for(; (i_34 < 1); i_34 = (i_34 + 1)) { - i_35 = 0; - { - [loop] for(; (i_35 < 1); i_35 = (i_35 + 1)) { - i_36 = 0; - { - [loop] for(; (i_36 < 1); i_36 = (i_36 + 1)) { - i_37 = 0; - { - [loop] for(; (i_37 < 1); i_37 = (i_37 + 1)) { - [loop] while (true) { - x_GLF_global_loop_count = (x_GLF_global_loop_count + 1); - { - if ((x_GLF_global_loop_count < 98)) { - } else { - break; - } - } - } - set_float3(m23[i_37], i_37, 1.0f); - set_float4(m24[i_37], i_37, 1.0f); - set_float2(m32[i_37], i_37, 1.0f); - set_float3(m33[i_37], i_37, 1.0f); - set_float4(m34[i_37], i_37, 1.0f); - set_float2(m42[i_37], i_37, 1.0f); - set_float3(m43[i_37], i_37, 1.0f); - set_float4(m44[i_37], i_37, 1.0f); - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - } - sum = 0.0f; - r = 0; - { - [loop] for(; (x_GLF_global_loop_count < 100); r = (r + 1)) { - x_GLF_global_loop_count = (x_GLF_global_loop_count + 1); - const float x_486 = m23[0][r]; - sum = (sum + x_486); - const float x_491 = m24[0][r]; - sum = (sum + x_491); - const float x_496 = m32[0][r]; - sum = (sum + x_496); - const float x_501 = m33[0][r]; - sum = (sum + x_501); - const float x_506 = m34[0][r]; - sum = (sum + x_506); - const float x_511 = m42[0][r]; - sum = (sum + x_511); - const float x_516 = m43[0][r]; - sum = (sum + x_516); - const float x_521 = m44[0][r]; - sum = (sum + x_521); - } - } - if ((sum == 8.0f)) { - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - } else { - x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_1 = {x_GLF_color}; - return tint_symbol_1; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x00000295E0FBF060(187,160-195): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.spvasm.expected.hlsl deleted file mode 100644 index 3b595c3804..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,97 +0,0 @@ -SKIP: FAILED - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_5 : register(b0, space0) { - uint4 x_5[4]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x2 m = float2x2(0.0f, 0.0f, 0.0f, 0.0f); - float f = 0.0f; - int i = 0; - int j = 0; - const int x_36 = asint(x_5[1].x); - if ((x_36 == 1)) { - const float x_40 = f; - m = float2x2(float2(x_40, 0.0f), float2(0.0f, x_40)); - } - const int x_45 = asint(x_5[1].x); - i = x_45; - [loop] while (true) { - const int x_50 = i; - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_52 = asint(x_5[scalar_offset / 4][scalar_offset % 4]); - if ((x_50 < x_52)) { - } else { - break; - } - const int x_56 = asint(x_5[1].x); - j = x_56; - [loop] while (true) { - const int x_61 = j; - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_63 = asint(x_5[scalar_offset_1 / 4][scalar_offset_1 % 4]); - if ((x_61 < x_63)) { - } else { - break; - } - const int x_66 = i; - const int x_67 = j; - const int x_68 = i; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const int x_70 = asint(x_5[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float2(m[x_66], x_67, float(((x_68 * x_70) + j))); - { - j = (j + 1); - } - } - { - i = (i + 1); - } - } - const float2x2 x_80 = m; - const int x_82 = asint(x_5[1].x); - const int x_85 = asint(x_5[2].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_88 = asint(x_5[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const int x_91 = asint(x_5[3].x); - const float2x2 x_95 = float2x2(float2(float(x_82), float(x_85)), float2(float(x_88), float(x_91))); - if ((all((x_80[0u] == x_95[0u])) & all((x_80[1u] == x_95[1u])))) { - const int x_109 = asint(x_5[2].x); - const int x_112 = asint(x_5[1].x); - const int x_115 = asint(x_5[1].x); - const int x_118 = asint(x_5[2].x); - x_GLF_color = float4(float(x_109), float(x_112), float(x_115), float(x_118)); - } else { - const int x_122 = asint(x_5[1].x); - const float x_123 = float(x_122); - x_GLF_color = float4(x_123, x_123, x_123, x_123); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_2 = {x_GLF_color}; - return tint_symbol_2; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000027B69224CC0(32,12-23): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 694f101ecd..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-reinitialize-matrix-after-undefined-value/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,101 +0,0 @@ -SKIP: FAILED - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_5 : register(b0, space0) { - uint4 x_5[4]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float2x2 m = float2x2(0.0f, 0.0f, 0.0f, 0.0f); - float f = 0.0f; - int i = 0; - int j = 0; - const int x_36 = asint(x_5[1].x); - if ((x_36 == 1)) { - const float x_40 = f; - m = float2x2(float2(x_40, 0.0f), float2(0.0f, x_40)); - } - const int x_45 = asint(x_5[1].x); - i = x_45; - [loop] while (true) { - const int x_50 = i; - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_52 = asint(x_5[scalar_offset / 4][scalar_offset % 4]); - if ((x_50 < x_52)) { - } else { - break; - } - const int x_56 = asint(x_5[1].x); - j = x_56; - [loop] while (true) { - const int x_61 = j; - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_63 = asint(x_5[scalar_offset_1 / 4][scalar_offset_1 % 4]); - if ((x_61 < x_63)) { - } else { - break; - } - const int x_66 = i; - const int x_67 = j; - const int x_68 = i; - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const int x_70 = asint(x_5[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float2(m[x_66], x_67, float(((x_68 * x_70) + j))); - { - j = (j + 1); - } - } - { - i = (i + 1); - } - } - const float2x2 x_80 = m; - const int x_82 = asint(x_5[1].x); - const int x_85 = asint(x_5[2].x); - const uint scalar_offset_3 = ((16u * uint(0))) / 4; - const int x_88 = asint(x_5[scalar_offset_3 / 4][scalar_offset_3 % 4]); - const int x_91 = asint(x_5[3].x); - const float2x2 x_95 = float2x2(float2(float(x_82), float(x_85)), float2(float(x_88), float(x_91))); - bool tint_tmp = all((x_80[0u] == x_95[0u])); - if (tint_tmp) { - tint_tmp = all((x_80[1u] == x_95[1u])); - } - if ((tint_tmp)) { - const int x_109 = asint(x_5[2].x); - const int x_112 = asint(x_5[1].x); - const int x_115 = asint(x_5[1].x); - const int x_118 = asint(x_5[2].x); - x_GLF_color = float4(float(x_109), float(x_112), float(x_115), float(x_118)); - } else { - const int x_122 = asint(x_5[1].x); - const float x_123 = float(x_122); - x_GLF_color = float4(x_123, x_123, x_123, x_123); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_2 = {x_GLF_color}; - return tint_symbol_2; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000019F4EA2D3B0(32,12-23): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.spvasm.expected.hlsl index 46469b7d05..5b9bdf5d61 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.spvasm.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.spvasm.expected.hlsl @@ -1,7 +1,20 @@ SKIP: FAILED -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; +void set_scalar_float4x3(inout float4x3 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[0]; + break; + case 1: + mat[1] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[1]; + break; + case 2: + mat[2] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[2]; + break; + case 3: + mat[3] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[3]; + break; + } } struct tint_padded_array_element { @@ -31,7 +44,7 @@ void main_1() { const int x_53 = asint(x_8[scalar_offset_1 / 4][scalar_offset_1 % 4]); const uint scalar_offset_2 = ((16u * uint(0))) / 4; const float x_55 = asfloat(x_6[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float3(m43[x_51], x_53, x_55); + set_scalar_float4x3(m43, x_53, x_51, x_55); const uint scalar_offset_3 = ((16u * uint(0))) / 4; const float x_58 = asfloat(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); const uint scalar_offset_4 = ((16u * uint(0))) / 4; @@ -115,5 +128,5 @@ tint_symbol main() { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x000001DF1D9D5EE0(32,14-22): error X3500: array reference cannot be used as an l-value; not natively addressable +C:\src\tint\test\Shader@0x00000111C3C21900(84,24-29): error X3504: array index out of bounds diff --git a/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.wgsl.expected.hlsl index 903d521787..5ab81af4ec 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.wgsl.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-unused-access-past-matrix-elements/0-opt.wgsl.expected.hlsl @@ -1,7 +1,20 @@ SKIP: FAILED -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; +void set_scalar_float4x3(inout float4x3 mat, int col, int row, float val) { + switch (col) { + case 0: + mat[0] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[0]; + break; + case 1: + mat[1] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[1]; + break; + case 2: + mat[2] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[2]; + break; + case 3: + mat[3] = (row.xxx == int3(0, 1, 2)) ? val.xxx : mat[3]; + break; + } } struct tint_padded_array_element { @@ -31,7 +44,7 @@ void main_1() { const int x_53 = asint(x_8[scalar_offset_1 / 4][scalar_offset_1 % 4]); const uint scalar_offset_2 = ((16u * uint(0))) / 4; const float x_55 = asfloat(x_6[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float3(m43[x_51], x_53, x_55); + set_scalar_float4x3(m43, x_53, x_51, x_55); const uint scalar_offset_3 = ((16u * uint(0))) / 4; const float x_58 = asfloat(x_6[scalar_offset_3 / 4][scalar_offset_3 % 4]); const uint scalar_offset_4 = ((16u * uint(0))) / 4; @@ -115,5 +128,5 @@ tint_symbol main() { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x0000020E5111BFA0(32,14-22): error X3500: array reference cannot be used as an l-value; not natively addressable +C:\src\tint\test\Shader@0x0000020DD7F92800(84,24-29): error X3504: array index out of bounds diff --git a/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.spvasm.expected.hlsl deleted file mode 100644 index 9f84295ee7..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,102 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -cbuffer cbuffer_x_6 : register(b1, space0) { - uint4 x_6[4]; -}; -cbuffer cbuffer_x_10 : register(b0, space0) { - uint4 x_10[1]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x4 m0 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int c = 0; - float4x4 m1 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const int x_40 = asint(x_6[1].x); - const float x_41 = float(x_40); - m0 = float4x4(float4(x_41, 0.0f, 0.0f, 0.0f), float4(0.0f, x_41, 0.0f, 0.0f), float4(0.0f, 0.0f, x_41, 0.0f), float4(0.0f, 0.0f, 0.0f, x_41)); - const int x_48 = asint(x_6[2].x); - c = x_48; - [loop] while (true) { - const int x_53 = c; - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_55 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - if ((x_53 < x_55)) { - } else { - break; - } - m1 = m0; - const int x_59 = c; - const int x_61 = asint(x_6[3].x); - const int x_64 = asint(x_6[2].x); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_66 = asfloat(x_10[scalar_offset_1 / 4][scalar_offset_1 % 4]); - set_float4(m1[(x_59 % x_61)], x_64, x_66); - const int x_68 = c; - const int x_70 = asint(x_6[3].x); - const int x_73 = asint(x_6[2].x); - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_75 = asfloat(x_10[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float4(m0[(x_68 % x_70)], x_73, x_75); - { - c = (c + 1); - } - } - const float4x4 x_79 = m0; - const int x_81 = asint(x_6[1].x); - const int x_84 = asint(x_6[2].x); - const int x_87 = asint(x_6[1].x); - const int x_90 = asint(x_6[1].x); - const int x_93 = asint(x_6[1].x); - const int x_96 = asint(x_6[2].x); - const int x_99 = asint(x_6[1].x); - const int x_102 = asint(x_6[1].x); - const int x_105 = asint(x_6[1].x); - const int x_108 = asint(x_6[2].x); - const int x_111 = asint(x_6[1].x); - const int x_114 = asint(x_6[1].x); - const int x_117 = asint(x_6[1].x); - const int x_120 = asint(x_6[2].x); - const int x_123 = asint(x_6[1].x); - const int x_126 = asint(x_6[1].x); - const float4x4 x_132 = float4x4(float4(float(x_81), float(x_84), float(x_87), float(x_90)), float4(float(x_93), float(x_96), float(x_99), float(x_102)), float4(float(x_105), float(x_108), float(x_111), float(x_114)), float4(float(x_117), float(x_120), float(x_123), float(x_126))); - if ((((all((x_79[0u] == x_132[0u])) & all((x_79[1u] == x_132[1u]))) & all((x_79[2u] == x_132[2u]))) & all((x_79[3u] == x_132[3u])))) { - const int x_156 = asint(x_6[2].x); - const int x_159 = asint(x_6[1].x); - const int x_162 = asint(x_6[1].x); - const int x_165 = asint(x_6[2].x); - x_GLF_color = float4(float(x_156), float(x_159), float(x_162), float(x_165)); - } else { - const int x_169 = asint(x_6[1].x); - const float x_170 = float(x_169); - x_GLF_color = float4(x_170, x_170, x_170, x_170); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_3 = {x_GLF_color}; - return tint_symbol_3; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000013C243CAFF0(36,20-30): warning X3556: integer modulus may be much slower, try using uints if possible. -C:\src\tint\test\Shader@0x0000013C243CAFF0(22,10-21): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.wgsl.expected.hlsl deleted file mode 100644 index 042c5abeb3..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-unused-matrix-copy-inside-loop/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,114 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -cbuffer cbuffer_x_6 : register(b1, space0) { - uint4 x_6[4]; -}; -cbuffer cbuffer_x_10 : register(b0, space0) { - uint4 x_10[1]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x4 m0 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int c = 0; - float4x4 m1 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - const int x_40 = asint(x_6[1].x); - const float x_41 = float(x_40); - m0 = float4x4(float4(x_41, 0.0f, 0.0f, 0.0f), float4(0.0f, x_41, 0.0f, 0.0f), float4(0.0f, 0.0f, x_41, 0.0f), float4(0.0f, 0.0f, 0.0f, x_41)); - const int x_48 = asint(x_6[2].x); - c = x_48; - [loop] while (true) { - const int x_53 = c; - const uint scalar_offset = ((16u * uint(0))) / 4; - const int x_55 = asint(x_6[scalar_offset / 4][scalar_offset % 4]); - if ((x_53 < x_55)) { - } else { - break; - } - m1 = m0; - const int x_59 = c; - const int x_61 = asint(x_6[3].x); - const int x_64 = asint(x_6[2].x); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const float x_66 = asfloat(x_10[scalar_offset_1 / 4][scalar_offset_1 % 4]); - set_float4(m1[(x_59 % x_61)], x_64, x_66); - const int x_68 = c; - const int x_70 = asint(x_6[3].x); - const int x_73 = asint(x_6[2].x); - const uint scalar_offset_2 = ((16u * uint(0))) / 4; - const float x_75 = asfloat(x_10[scalar_offset_2 / 4][scalar_offset_2 % 4]); - set_float4(m0[(x_68 % x_70)], x_73, x_75); - { - c = (c + 1); - } - } - const float4x4 x_79 = m0; - const int x_81 = asint(x_6[1].x); - const int x_84 = asint(x_6[2].x); - const int x_87 = asint(x_6[1].x); - const int x_90 = asint(x_6[1].x); - const int x_93 = asint(x_6[1].x); - const int x_96 = asint(x_6[2].x); - const int x_99 = asint(x_6[1].x); - const int x_102 = asint(x_6[1].x); - const int x_105 = asint(x_6[1].x); - const int x_108 = asint(x_6[2].x); - const int x_111 = asint(x_6[1].x); - const int x_114 = asint(x_6[1].x); - const int x_117 = asint(x_6[1].x); - const int x_120 = asint(x_6[2].x); - const int x_123 = asint(x_6[1].x); - const int x_126 = asint(x_6[1].x); - const float4x4 x_132 = float4x4(float4(float(x_81), float(x_84), float(x_87), float(x_90)), float4(float(x_93), float(x_96), float(x_99), float(x_102)), float4(float(x_105), float(x_108), float(x_111), float(x_114)), float4(float(x_117), float(x_120), float(x_123), float(x_126))); - bool tint_tmp_2 = all((x_79[0u] == x_132[0u])); - if (tint_tmp_2) { - tint_tmp_2 = all((x_79[1u] == x_132[1u])); - } - bool tint_tmp_1 = (tint_tmp_2); - if (tint_tmp_1) { - tint_tmp_1 = all((x_79[2u] == x_132[2u])); - } - bool tint_tmp = (tint_tmp_1); - if (tint_tmp) { - tint_tmp = all((x_79[3u] == x_132[3u])); - } - if ((tint_tmp)) { - const int x_156 = asint(x_6[2].x); - const int x_159 = asint(x_6[1].x); - const int x_162 = asint(x_6[1].x); - const int x_165 = asint(x_6[2].x); - x_GLF_color = float4(float(x_156), float(x_159), float(x_162), float(x_165)); - } else { - const int x_169 = asint(x_6[1].x); - const float x_170 = float(x_169); - x_GLF_color = float4(x_170, x_170, x_170, x_170); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_3 = {x_GLF_color}; - return tint_symbol_3; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001ABB4809E40(36,20-30): warning X3556: integer modulus may be much slower, try using uints if possible. -C:\src\tint\test\Shader@0x000001ABB4809E40(22,10-21): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.spvasm.expected.hlsl deleted file mode 100755 index 6760777033..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.spvasm.expected.hlsl +++ /dev/null @@ -1,75 +0,0 @@ -SKIP: FAILED - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_6 : register(b1, space0) { - uint4 x_6[2]; -}; -cbuffer cbuffer_x_8 : register(b0, space0) { - uint4 x_8[3]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float3x2 m32 = float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float sums[3] = (float[3])0; - int x_52_phi = 0; - const uint scalar_offset = ((16u * uint(0))) / 4; - const float x_40 = asfloat(x_6[scalar_offset / 4][scalar_offset % 4]); - m32 = float3x2(float2(x_40, 0.0f), float2(0.0f, x_40), float2(0.0f, 0.0f)); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_45 = asint(x_8[scalar_offset_1 / 4][scalar_offset_1 % 4]); - if ((x_45 == 1)) { - set_float2(m32[3], x_45, x_40); - } - const float tint_symbol_3[3] = {x_40, x_40, x_40}; - sums = tint_symbol_3; - x_52_phi = x_45; - [loop] while (true) { - int x_53 = 0; - const int x_52 = x_52_phi; - const int x_56 = asint(x_8[2].x); - if ((x_52 < x_56)) { - } else { - break; - } - { - const float x_60 = m32[x_52][x_45]; - const int x_61_save = x_56; - const float x_62 = sums[x_61_save]; - sums[x_61_save] = (x_62 + x_60); - x_53 = (x_52 + 1); - x_52_phi = x_53; - } - } - const float x_65 = sums[x_45]; - const float x_67 = asfloat(x_6[1].x); - const int x_69 = asint(x_8[1].x); - const float x_71 = sums[x_69]; - x_GLF_color = float4(x_65, x_67, x_67, x_71); - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000018CEDCE0230(23,16-21): error X3504: array index out of bounds - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.wgsl.expected.hlsl deleted file mode 100755 index c4c7d329b0..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/cov-write-past-matrix-elements-unused/0.wgsl.expected.hlsl +++ /dev/null @@ -1,75 +0,0 @@ -SKIP: FAILED - -void set_float2(inout float2 vec, int idx, float val) { - vec = (idx.xx == int2(0, 1)) ? val.xx : vec; -} - -cbuffer cbuffer_x_6 : register(b1, space0) { - uint4 x_6[2]; -}; -cbuffer cbuffer_x_8 : register(b0, space0) { - uint4 x_8[3]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float3x2 m32 = float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - float sums[3] = (float[3])0; - int x_52_phi = 0; - const uint scalar_offset = ((16u * uint(0))) / 4; - const float x_40 = asfloat(x_6[scalar_offset / 4][scalar_offset % 4]); - m32 = float3x2(float2(x_40, 0.0f), float2(0.0f, x_40), float2(0.0f, 0.0f)); - const uint scalar_offset_1 = ((16u * uint(0))) / 4; - const int x_45 = asint(x_8[scalar_offset_1 / 4][scalar_offset_1 % 4]); - if ((x_45 == 1)) { - set_float2(m32[3], x_45, x_40); - } - const float tint_symbol_3[3] = {x_40, x_40, x_40}; - sums = tint_symbol_3; - x_52_phi = x_45; - [loop] while (true) { - int x_53 = 0; - const int x_52 = x_52_phi; - const int x_56 = asint(x_8[2].x); - if ((x_52 < x_56)) { - } else { - break; - } - { - const float x_60 = m32[x_52][x_45]; - const int x_61_save = x_56; - const float x_62 = sums[x_61_save]; - sums[x_61_save] = (x_62 + x_60); - x_53 = (x_52 + 1); - x_52_phi = x_53; - } - } - const float x_65 = sums[x_45]; - const float x_67 = asfloat(x_6[1].x); - const int x_69 = asint(x_8[1].x); - const float x_71 = sums[x_69]; - x_GLF_color = float4(x_65, x_67, x_67, x_71); - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x00000229BC652210(23,16-21): error X3504: array index out of bounds - diff --git a/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.spvasm.expected.hlsl deleted file mode 100644 index dea6d2c36a..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.spvasm.expected.hlsl +++ /dev/null @@ -1,86 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -cbuffer cbuffer_x_7 : register(b0, space0) { - uint4 x_7[1]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x4 m44 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int x_10_phi = 0; - m44 = float4x4(float4(1.0f, 2.0f, 3.0f, 4.0f), float4(5.0f, 6.0f, 7.0f, 8.0f), float4(9.0f, 10.0f, 11.0f, 12.0f), float4(13.0f, 14.0f, 15.0f, 16.0f)); - x_10_phi = 0; - [loop] while (true) { - int x_9 = 0; - int x_11_phi = 0; - const int x_10 = x_10_phi; - if ((x_10 < 4)) { - } else { - break; - } - const float x_63 = gl_FragCoord.y; - if ((x_63 < 0.0f)) { - break; - } - x_11_phi = 0; - [loop] while (true) { - int x_8 = 0; - const int x_11 = x_11_phi; - if ((x_11 < 4)) { - } else { - break; - } - { - const float x_72 = asfloat(x_7[0].x); - const float x_74 = m44[x_10][x_11]; - set_float4(m44[x_10], x_11, (x_74 + x_72)); - x_8 = (x_11 + 1); - x_11_phi = x_8; - } - } - { - x_9 = (x_10 + 1); - x_10_phi = x_9; - } - } - const float x_77 = m44[1].y; - float4 x_79_1 = float4(0.0f, 0.0f, 0.0f, 0.0f); - x_79_1.x = (x_77 - 6.0f); - const float4 x_79 = x_79_1; - const float x_81 = m44[2].z; - float4 x_83_1 = x_79; - x_83_1.w = (x_81 - 11.0f); - x_GLF_color = x_83_1; - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001CE5E694FB0(29,12-23): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.wgsl.expected.hlsl deleted file mode 100644 index 59429ad413..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/nested-for-break-mat-color/0.wgsl.expected.hlsl +++ /dev/null @@ -1,86 +0,0 @@ -SKIP: FAILED - -void set_float4(inout float4 vec, int idx, float val) { - vec = (idx.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : vec; -} - -static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); -cbuffer cbuffer_x_7 : register(b0, space0) { - uint4 x_7[1]; -}; -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x4 m44 = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int x_10_phi = 0; - m44 = float4x4(float4(1.0f, 2.0f, 3.0f, 4.0f), float4(5.0f, 6.0f, 7.0f, 8.0f), float4(9.0f, 10.0f, 11.0f, 12.0f), float4(13.0f, 14.0f, 15.0f, 16.0f)); - x_10_phi = 0; - [loop] while (true) { - int x_9 = 0; - int x_11_phi = 0; - const int x_10 = x_10_phi; - if ((x_10 < 4)) { - } else { - break; - } - const float x_63 = gl_FragCoord.y; - if ((x_63 < 0.0f)) { - break; - } - x_11_phi = 0; - [loop] while (true) { - int x_8 = 0; - const int x_11 = x_11_phi; - if ((x_11 < 4)) { - } else { - break; - } - { - const float x_72 = asfloat(x_7[0].x); - const float x_74 = m44[x_10][x_11]; - set_float4(m44[x_10], x_11, (x_74 + x_72)); - x_8 = (x_11 + 1); - x_11_phi = x_8; - } - } - { - x_9 = (x_10 + 1); - x_10_phi = x_9; - } - } - const float x_77 = m44[1].y; - float4 x_79_1 = float4(0.0f, 0.0f, 0.0f, 0.0f); - x_79_1.x = (x_77 - 6.0f); - const float4 x_79 = x_79_1; - const float x_81 = m44[2].z; - float4 x_83_1 = x_79; - x_83_1.w = (x_81 - 11.0f); - x_GLF_color = x_83_1; - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol_1 { - float4 gl_FragCoord_param : SV_Position; -}; -struct tint_symbol_2 { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner(float4 gl_FragCoord_param) { - gl_FragCoord = gl_FragCoord_param; - main_1(); - const main_out tint_symbol_4 = {x_GLF_color}; - return tint_symbol_4; -} - -tint_symbol_2 main(tint_symbol_1 tint_symbol) { - const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param); - tint_symbol_2 wrapper_result = (tint_symbol_2)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001A0A52BEBB0(29,12-23): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.spvasm.expected.hlsl deleted file mode 100644 index 8a1a43ab85..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.spvasm.expected.hlsl +++ /dev/null @@ -1,83 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x3 m43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int ll1 = 0; - int rows = 0; - int ll4 = 0; - int ll2 = 0; - int c = 0; - float4x3 tempm43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int ll3 = 0; - int d = 0; - int r = 0; - float sums[9] = (float[9])0; - int idx = 0; - m43 = float4x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f), float3(0.0f, 0.0f, 0.0f)); - ll1 = 0; - rows = 2; - [loop] while (true) { - if (true) { - } else { - break; - } - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - if ((ll1 >= 5)) { - break; - } - ll1 = (ll1 + 1); - ll4 = 10; - ll2 = 0; - c = 0; - { - [loop] for(; (c < 1); c = (c + 1)) { - if ((ll2 >= 0)) { - break; - } - ll2 = (ll2 + 1); - tempm43 = m43; - ll3 = 0; - d = 0; - { - [loop] for(; (1 < ll4); d = (d + 1)) { - set_float3(tempm43[(((d >= 0) & (d < 4)) ? d : 0)], (((r >= 0) & (r < 3)) ? r : 0), 1.0f); - } - } - const int x_111 = (((idx >= 0) & (idx < 9)) ? idx : 0); - const float x_113 = m43[c].y; - const float x_115 = sums[x_111]; - sums[x_111] = (x_115 + x_113); - } - } - idx = (idx + 1); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_1 = {x_GLF_color}; - return tint_symbol_1; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x000001DA1CD04F80(46,18-46): error X3531: can't unroll loops marked with loop attribute - diff --git a/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.wgsl.expected.hlsl deleted file mode 100644 index caf9354e1c..0000000000 --- a/test/vk-gl-cts/graphicsfuzz/write-red-in-loop-nest/0-opt.wgsl.expected.hlsl +++ /dev/null @@ -1,95 +0,0 @@ -SKIP: FAILED - -void set_float3(inout float3 vec, int idx, float val) { - vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec; -} - -static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); - -void main_1() { - float4x3 m43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int ll1 = 0; - int rows = 0; - int ll4 = 0; - int ll2 = 0; - int c = 0; - float4x3 tempm43 = float4x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); - int ll3 = 0; - int d = 0; - int r = 0; - float sums[9] = (float[9])0; - int idx = 0; - m43 = float4x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f), float3(0.0f, 0.0f, 0.0f)); - ll1 = 0; - rows = 2; - [loop] while (true) { - if (true) { - } else { - break; - } - x_GLF_color = float4(1.0f, 0.0f, 0.0f, 1.0f); - if ((ll1 >= 5)) { - break; - } - ll1 = (ll1 + 1); - ll4 = 10; - ll2 = 0; - c = 0; - { - [loop] for(; (c < 1); c = (c + 1)) { - if ((ll2 >= 0)) { - break; - } - ll2 = (ll2 + 1); - tempm43 = m43; - ll3 = 0; - d = 0; - { - [loop] for(; (1 < ll4); d = (d + 1)) { - bool tint_tmp = (d >= 0); - if (tint_tmp) { - tint_tmp = (d < 4); - } - bool tint_tmp_1 = (r >= 0); - if (tint_tmp_1) { - tint_tmp_1 = (r < 3); - } - set_float3(tempm43[((tint_tmp) ? d : 0)], ((tint_tmp_1) ? r : 0), 1.0f); - } - } - bool tint_tmp_2 = (idx >= 0); - if (tint_tmp_2) { - tint_tmp_2 = (idx < 9); - } - const int x_111 = ((tint_tmp_2) ? idx : 0); - const float x_113 = m43[c].y; - const float x_115 = sums[x_111]; - sums[x_111] = (x_115 + x_113); - } - } - idx = (idx + 1); - } - return; -} - -struct main_out { - float4 x_GLF_color_1; -}; -struct tint_symbol { - float4 x_GLF_color_1 : SV_Target0; -}; - -main_out main_inner() { - main_1(); - const main_out tint_symbol_1 = {x_GLF_color}; - return tint_symbol_1; -} - -tint_symbol main() { - const main_out inner_result = main_inner(); - tint_symbol wrapper_result = (tint_symbol)0; - wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; - return wrapper_result; -} -C:\src\tint\test\Shader@0x0000026405D97AD0(46,18-46): error X3531: can't unroll loops marked with loop attribute -