HLSL: work around FXC failures when dynamically indexing matrices
This fixes errors like "error X3500: array reference cannot be used as
an l-value; not natively addressable". Note that FXC treats matrices
like arrays. We still get this error for dynamically indexed arrays in
structs.
Also improved HLSL assign tests, and add missing ones for vector
indexing.
Manually removed 20 e2e skip hlsl SKIP files that are now passing with
this change.
Bug: tint:1333
Change-Id: If23881a667857a4d4ec6881e72666af0a666ef10
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/71982
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 122f50e..c66278f 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -344,6 +344,181 @@
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<ast::IndexAccessorExpression>();
+
+ 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<ast::IndexAccessorExpression>();
+ auto* lhs_row_access =
+ lhs_col_access->object->As<ast::IndexAccessorExpression>();
+
+ 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<sem::Vector>();
+ 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::EmitAssign(const ast::AssignmentStatement* stmt) {
- if (auto* idx = stmt->lhs->As<ast::IndexAccessorExpression>()) {
- if (auto* vec = TypeOf(idx->object)->UnwrapRef()->As<sem::Vector>()) {
- auto* rhs_sem = builder_.Sem().Get(idx->index);
+ if (auto* lhs_access = stmt->lhs->As<ast::IndexAccessorExpression>()) {
+ // 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<ast::IndexAccessorExpression>()) {
+ if (auto* mat =
+ TypeOf(lhs_sub_access->object)->UnwrapRef()->As<sem::Matrix>()) {
+ 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<sem::Matrix>()) {
+ 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<sem::Vector>()) {
+ 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 c2532a2..12124d2 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -423,6 +423,26 @@
/// @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 @@
std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
std::unordered_map<const sem::Struct*, std::string> structure_builders_;
std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
+ std::unordered_map<const sem::Matrix*, std::string>
+ dynamic_matrix_vector_write_;
+ std::unordered_map<const sem::Matrix*, std::string>
+ 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 f5a0711..98d9eb6 100644
--- a/src/writer/hlsl/generator_impl_assign_test.cc
+++ b/src/writer/hlsl/generator_impl_assign_test.cc
@@ -22,17 +22,187 @@
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<f32>())),
+ 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<f32>())),
+ 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<f32>())),
+ Decl(Var("rhs", ty.vec2<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() {
+ 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<f32>())),
+ Decl(Var("rhs", ty.vec2<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_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<f32>())),
+ 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<f32>())),
+ 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 0000000..7c39426
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ 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 0000000..d9c9833
--- /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 0000000..82056b4
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..c2b242a
--- /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 0000000..eeaa5a9
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ 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 0000000..9f5e961
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ 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 0000000..254519f
--- /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 0000000..ea2b9ba
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..b3c73b4
--- /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 0000000..b4c8fe9
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ 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 0000000..345e6e0
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..75a6217
--- /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 0000000..4e4bb57
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..3f68ce1
--- /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 0000000..9b78e5a4
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..fa0ad25
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ m1[uniforms.i] = vec4<f32>(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 0000000..4b5af71
--- /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 0000000..a5aed7f
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..afc1c44
--- /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 0000000..ef9e7ad
--- /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<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var m1 : mat2x4<f32>;
+ m1[uniforms.i] = vec4<f32>(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 0000000..5325331
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..c52f157
--- /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 0000000..31c3003
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..b2bb365
--- /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 0000000..ea6fc1b
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..9ccb54a
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..1487ddd
--- /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 0000000..7215cb0
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..1b93b11
--- /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 0000000..b9d60a2
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..345e6e0
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..75a6217
--- /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 0000000..4e4bb57
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..3f68ce1
--- /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 0000000..9b78e5a4
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[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 0000000..d5b3c63
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ m1[uniforms.i] = vec4<f32>(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 0000000..bf03373
--- /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 0000000..6c3bb03
--- /dev/null
+++ b/test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+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 0000000..f669264
--- /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 0000000..acafd9d
--- /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<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ m1[uniforms.i] = vec4<f32>(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 f78f210..0000000
--- 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 6b45ab8..0000000
--- 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 909cf6f..0000000
--- 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 83a7b7d..0000000
--- 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 e2a6cb3..0000000
--- 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 884c3d6..0000000
--- 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 097dd84..0000000
--- 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 eb7806c..0000000
--- 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 cfed15c..0000000
--- 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 1b87903..0000000
--- 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 3b595c3..0000000
--- 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 694f101..0000000
--- 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 46469b7..5b9bdf5 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 @@
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 @@
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 903d521..5ab81af 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 @@
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 @@
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 9f84295..0000000
--- 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 042c5ab..0000000
--- 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 6760777..0000000
--- 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 c4c7d32..0000000
--- 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 dea6d2c..0000000
--- 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 59429ad..0000000
--- 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 8a1a43a..0000000
--- 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 caf9354..0000000
--- 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
-
