tint: Flip evaluation order of assignment statements
Evaluate the LHS before the RHS.
Fixed: tint:1867
Change-Id: Ib63903ed4b1425007197a6da37f3bf54a495d88a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/123120
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc
index bc4ac1d..673c1c0 100644
--- a/src/tint/resolver/uniformity.cc
+++ b/src/tint/resolver/uniformity.cc
@@ -45,6 +45,10 @@
// Set to `1` to dump the uniformity graph for each function in graphviz format.
#define TINT_DUMP_UNIFORMITY_GRAPH 0
+#if TINT_DUMP_UNIFORMITY_GRAPH
+#include <iostream>
+#endif
+
namespace tint::resolver {
namespace {
@@ -123,7 +127,7 @@
const ast::Node* ast = nullptr;
/// The function call argument index, if applicable.
- uint32_t arg_index;
+ uint32_t arg_index = 0xffffffffu;
/// The set of edges from this node to other nodes in the graph.
utils::UniqueVector<Node*, 4> edges;
@@ -547,14 +551,15 @@
stmt,
[&](const ast::AssignmentStatement* a) {
- auto [cf1, v1] = ProcessExpression(cf, a->rhs);
if (a->lhs->Is<ast::PhonyExpression>()) {
- return cf1;
- } else {
- auto [cf2, l2] = ProcessLValueExpression(cf1, a->lhs);
- l2->AddEdge(v1);
- return cf2;
+ auto [cf_r, _] = ProcessExpression(cf, a->rhs);
+ return cf_r;
}
+ auto [cf_l, v_l, apply] = ProcessLValueExpression(cf, a->lhs);
+ auto [cf_r, v_r] = ProcessExpression(cf_l, a->rhs);
+ v_l->AddEdge(v_r);
+ apply();
+ return cf_r;
},
[&](const ast::BlockStatement* b) {
@@ -696,17 +701,20 @@
},
[&](const ast::CompoundAssignmentStatement* c) {
- // The compound assignment statement `a += b` is equivalent to `a = a + b`.
- // Note: we set load_rule=true when evaluating the LHS the first time, as the
- // resolver does not add a load node for it.
- auto [cf1, v1] = ProcessExpression(cf, c->lhs, /* load_rule */ true);
- auto [cf2, v2] = ProcessExpression(cf1, c->rhs);
+ // The compound assignment statement `a += b` is equivalent to:
+ // let p = &a;
+ // *p = *p + b;
+ // Note: we set load_rule=true when evaluating the LHS, as the resolver does not add
+ // a load node for it.
+ auto [cf1, l1, apply] = ProcessLValueExpression(cf, c->lhs);
+ auto [cf2, v2] = ProcessExpression(cf1, c->lhs, /* load_rule */ true);
+ auto [cf3, v3] = ProcessExpression(cf2, c->rhs);
auto* result = CreateNode({"binary_expr_result"});
- result->AddEdge(v1);
result->AddEdge(v2);
+ result->AddEdge(v3);
- auto [cf3, l3] = ProcessLValueExpression(cf2, c->lhs);
- l3->AddEdge(result);
+ l1->AddEdge(result);
+ apply();
return cf3;
},
@@ -965,8 +973,9 @@
result->AddEdge(v1);
result->AddEdge(cf1);
- auto [cf2, l2] = ProcessLValueExpression(cf1, i->lhs);
+ auto [cf2, l2, apply] = ProcessLValueExpression(cf1, i->lhs);
l2->AddEdge(result);
+ apply();
return cf2;
},
@@ -1365,48 +1374,62 @@
return false;
}
+ /// LValue holds the Nodes returned by ProcessLValueExpression()
+ struct LValue {
+ /// The control-flow node for an LValue expression
+ Node* cf = nullptr;
+
+ /// The new value node for an LValue expression
+ Node* new_val = nullptr;
+
+ /// Updates the value node of the LValue expression to be #new_val.
+ std::function<void()> apply;
+ };
+
/// Process an LValue expression.
/// @param cf the input control flow node
/// @param expr the expression to process
/// @returns a pair of (control flow node, variable node)
- std::pair<Node*, Node*> ProcessLValueExpression(Node* cf,
- const ast::Expression* expr,
- bool is_partial_reference = false) {
+ LValue ProcessLValueExpression(Node* cf,
+ const ast::Expression* expr,
+ bool is_partial_reference = false) {
return Switch(
expr,
[&](const ast::IdentifierExpression* i) {
auto* sem = sem_.GetVal(i)->UnwrapLoad()->As<sem::VariableUser>();
if (sem->Variable()->Is<sem::GlobalVariable>()) {
- return std::make_pair(cf, current_function_->may_be_non_uniform);
+ return LValue{cf, current_function_->may_be_non_uniform, [] {}};
} else if (auto* local = sem->Variable()->As<sem::LocalVariable>()) {
// Create a new value node for this variable.
auto* value = CreateNode({NameFor(i), "_lvalue"});
- auto* old_value = current_function_->variables.Set(local, value);
+
+ auto apply = [=] { current_function_->variables.Set(local, value); };
// If i is part of an expression that is a partial reference to a variable (e.g.
// index or member access), we link back to the variable's previous value. If
// the previous value was non-uniform, a partial assignment will not make it
// uniform.
+ auto* old_value = current_function_->variables.Get(local);
if (is_partial_reference && old_value) {
value->AddEdge(old_value);
}
- return std::make_pair(cf, value);
+ return LValue{cf, value, apply};
} else {
TINT_ICE(Resolver, diagnostics_)
<< "unknown lvalue identifier expression type: "
<< std::string(sem->Variable()->TypeInfo().name);
- return std::pair<Node*, Node*>(nullptr, nullptr);
+ return LValue{};
}
},
[&](const ast::IndexAccessorExpression* i) {
- auto [cf1, l1] =
+ auto [cf1, l1, apply] =
ProcessLValueExpression(cf, i->object, /*is_partial_reference*/ true);
auto [cf2, v2] = ProcessExpression(cf1, i->index);
l1->AddEdge(v2);
- return std::pair<Node*, Node*>(cf2, l1);
+ return LValue{cf2, l1, apply};
},
[&](const ast::MemberAccessorExpression* m) {
@@ -1419,17 +1442,18 @@
// that is being written to.
auto* root_ident = sem_.Get(u)->RootIdentifier();
auto* deref = CreateNode({NameFor(root_ident), "_deref"});
- auto* old_value = current_function_->variables.Set(root_ident, deref);
- if (old_value) {
- // If derefercing a partial reference or partial pointer, we link back to
+ auto apply = [=] { current_function_->variables.Set(root_ident, deref); };
+
+ if (auto* old_value = current_function_->variables.Get(root_ident)) {
+ // If dereferencing a partial reference or partial pointer, we link back to
// the variable's previous value. If the previous value was non-uniform, a
// partial assignment will not make it uniform.
if (is_partial_reference || IsDerefOfPartialPointer(u)) {
deref->AddEdge(old_value);
}
}
- return std::pair<Node*, Node*>(cf, deref);
+ return LValue{cf, deref, apply};
}
return ProcessLValueExpression(cf, u->expr, is_partial_reference);
},
@@ -1437,7 +1461,7 @@
[&](Default) {
TINT_ICE(Resolver, diagnostics_)
<< "unknown lvalue expression type: " << std::string(expr->TypeInfo().name);
- return std::pair<Node*, Node*>(nullptr, nullptr);
+ return LValue{};
});
}
diff --git a/src/tint/resolver/uniformity_test.cc b/src/tint/resolver/uniformity_test.cc
index 51608bd..a55d259 100644
--- a/src/tint/resolver/uniformity_test.cc
+++ b/src/tint/resolver/uniformity_test.cc
@@ -8503,5 +8503,151 @@
)");
}
+TEST_F(UniformityAnalysisTest, AssignmentEval_LHS_Then_RHS_Pass) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn b(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return 0;
+}
+
+fn a(p : ptr<function, i32>) -> i32 {
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+ return 0;
+}
+
+fn foo() {
+ var i = 0;
+ var arr : array<i32, 4>;
+ arr[a(&i)] = arr[b(&i)];
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignmentEval_LHS_Then_RHS_Fail) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn a(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return 0;
+}
+
+fn b(p : ptr<function, i32>) -> i32 {
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+ return 0;
+}
+
+fn foo() {
+ var i = 0;
+ var arr : array<i32, 4>;
+ arr[a(&i)] = arr[b(&i)];
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:10:3 note: control flow depends on possibly non-uniform value
+ if (*p == 0) {
+ ^^
+
+test:10:8 note: parameter 'p' of 'b' may be non-uniform
+ if (*p == 0) {
+ ^
+
+test:19:22 note: possibly non-uniform value passed via pointer here
+ arr[a(&i)] = arr[b(&i)];
+ ^
+
+test:19:9 note: contents of pointer may become non-uniform after calling 'a'
+ arr[a(&i)] = arr[b(&i)];
+ ^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, CompoundAssignmentEval_LHS_Then_RHS_Pass) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn b(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return 0;
+}
+
+fn a(p : ptr<function, i32>) -> i32 {
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+ return 0;
+}
+
+fn foo() {
+ var i = 0;
+ var arr : array<i32, 4>;
+ arr[a(&i)] += arr[b(&i)];
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, CompoundAssignmentEval_LHS_Then_RHS_Fail) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn a(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return 0;
+}
+
+fn b(p : ptr<function, i32>) -> i32 {
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+ return 0;
+}
+
+fn foo() {
+ var i = 0;
+ var arr : array<i32, 4>;
+ arr[a(&i)] += arr[b(&i)];
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:10:3 note: control flow depends on possibly non-uniform value
+ if (*p == 0) {
+ ^^
+
+test:10:8 note: parameter 'p' of 'b' may be non-uniform
+ if (*p == 0) {
+ ^
+
+test:19:23 note: possibly non-uniform value passed via pointer here
+ arr[a(&i)] += arr[b(&i)];
+ ^
+
+test:19:9 note: contents of pointer may become non-uniform after calling 'a'
+ arr[a(&i)] += arr[b(&i)];
+ ^
+)");
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/transform/promote_side_effects_to_decl.cc b/src/tint/transform/promote_side_effects_to_decl.cc
index 65ec731..9db7e38 100644
--- a/src/tint/transform/promote_side_effects_to_decl.cc
+++ b/src/tint/transform/promote_side_effects_to_decl.cc
@@ -337,9 +337,8 @@
});
}
- // Starts the recursive processing of a statement's expression(s) to hoist
- // side-effects to lets.
- void ProcessStatement(const ast::Expression* expr) {
+ // Starts the recursive processing of a statement's expression(s) to hoist side-effects to lets.
+ void ProcessExpression(const ast::Expression* expr) {
if (!expr) {
return;
}
@@ -348,31 +347,6 @@
ProcessExpression(expr, maybe_hoist);
}
- // Special case for processing assignment statement expressions, as we must
- // evaluate the rhs before the lhs, and possibly hoist the rhs expression.
- void ProcessAssignment(const ast::Expression* lhs, const ast::Expression* rhs) {
- // Evaluate rhs before lhs
- tint::utils::Vector<const ast::Expression*, 8> maybe_hoist;
- if (ProcessExpression(rhs, maybe_hoist)) {
- maybe_hoist.Push(rhs);
- }
-
- // If the rhs has side-effects, it may affect the lhs, so hoist it right
- // away. e.g. "b[c] = a(0);"
- if (HasSideEffects(rhs)) {
- // Technically, we can always hoist rhs, but don't bother doing so when
- // the lhs is just a variable or phony.
- if (!lhs->IsAnyOf<ast::IdentifierExpression, ast::PhonyExpression>()) {
- Flush(maybe_hoist);
- }
- }
-
- // If maybe_hoist still has values, it means they are potential side-effect
- // receivers. We pass this in while processing the lhs, in which case they
- // may get hoisted if the lhs has side-effects. E.g. "b[a(0)] = c;".
- ProcessExpression(lhs, maybe_hoist);
- }
-
public:
explicit CollectHoistsState(CloneContext& ctx_in) : StateBase(ctx_in) {}
@@ -386,21 +360,26 @@
}
Switch(
- stmt, [&](const ast::AssignmentStatement* s) { ProcessAssignment(s->lhs, s->rhs); },
- [&](const ast::CallStatement* s) { //
- ProcessStatement(s->expr);
+ stmt, //
+ [&](const ast::AssignmentStatement* s) {
+ tint::utils::Vector<const ast::Expression*, 8> maybe_hoist;
+ ProcessExpression(s->lhs, maybe_hoist);
+ ProcessExpression(s->rhs, maybe_hoist);
},
- [&](const ast::ForLoopStatement* s) { ProcessStatement(s->condition); },
- [&](const ast::WhileStatement* s) { ProcessStatement(s->condition); },
+ [&](const ast::CallStatement* s) { //
+ ProcessExpression(s->expr);
+ },
+ [&](const ast::ForLoopStatement* s) { ProcessExpression(s->condition); },
+ [&](const ast::WhileStatement* s) { ProcessExpression(s->condition); },
[&](const ast::IfStatement* s) { //
- ProcessStatement(s->condition);
+ ProcessExpression(s->condition);
},
[&](const ast::ReturnStatement* s) { //
- ProcessStatement(s->value);
+ ProcessExpression(s->value);
},
- [&](const ast::SwitchStatement* s) { ProcessStatement(s->condition); },
+ [&](const ast::SwitchStatement* s) { ProcessExpression(s->condition); },
[&](const ast::VariableDeclStatement* s) {
- ProcessStatement(s->variable->initializer);
+ ProcessExpression(s->variable->initializer);
});
}
@@ -563,10 +542,10 @@
!sem.GetVal(s->rhs)->HasSideEffects()) {
return nullptr;
}
- // rhs before lhs
+ // lhs before rhs
tint::utils::Vector<const ast::Statement*, 8> stmts;
- ctx.Replace(s->rhs, Decompose(s->rhs, &stmts));
ctx.Replace(s->lhs, Decompose(s->lhs, &stmts));
+ ctx.Replace(s->rhs, Decompose(s->rhs, &stmts));
InsertBefore(stmts, s);
return ctx.CloneWithoutTransform(s);
},
diff --git a/src/tint/transform/promote_side_effects_to_decl_test.cc b/src/tint/transform/promote_side_effects_to_decl_test.cc
index 6ed4b27..636a809 100644
--- a/src/tint/transform/promote_side_effects_to_decl_test.cc
+++ b/src/tint/transform/promote_side_effects_to_decl_test.cc
@@ -2830,9 +2830,8 @@
fn f() {
var b = array<i32, 10>();
- let tint_symbol = a(1);
- let tint_symbol_1 = a(0);
- b[tint_symbol_1] = tint_symbol;
+ let tint_symbol = a(0);
+ b[tint_symbol] = a(1);
}
)";
@@ -2861,10 +2860,9 @@
fn f() {
var b = array<array<i32, 10>, 10>();
- let tint_symbol = a(2);
- let tint_symbol_1 = a(0);
- let tint_symbol_2 = a(1);
- b[tint_symbol_1][tint_symbol_2] = tint_symbol;
+ let tint_symbol = a(0);
+ let tint_symbol_1 = a(1);
+ b[tint_symbol][tint_symbol_1] = a(2);
}
)";
@@ -2893,11 +2891,10 @@
fn f() {
var b = array<array<array<i32, 10>, 10>, 10>();
- let tint_symbol = a(3);
- let tint_symbol_1 = a(0);
- let tint_symbol_2 = a(1);
- let tint_symbol_3 = a(2);
- b[tint_symbol_1][tint_symbol_2][tint_symbol_3] = tint_symbol;
+ let tint_symbol = a(0);
+ let tint_symbol_1 = a(1);
+ let tint_symbol_2 = a(2);
+ b[tint_symbol][tint_symbol_1][tint_symbol_2] = a(3);
}
)";
@@ -2930,11 +2927,9 @@
var b = array<i32, 3>();
var d = array<array<i32, 3>, 3>();
var a_1 = 0;
- let tint_symbol = a(0);
- let tint_symbol_1 = a_1;
- let tint_symbol_2 = d[tint_symbol][tint_symbol_1];
- let tint_symbol_3 = a(2);
- b[tint_symbol_3] = tint_symbol_2;
+ let tint_symbol = a(2);
+ let tint_symbol_1 = a(0);
+ b[tint_symbol] = d[tint_symbol_1][a_1];
}
)";
@@ -2963,9 +2958,8 @@
fn f() {
var b = vec3<i32>();
- let tint_symbol = a(1);
- let tint_symbol_1 = a(0);
- b[tint_symbol_1] = tint_symbol;
+ let tint_symbol = a(0);
+ b[tint_symbol] = a(1);
}
)";
@@ -2996,9 +2990,8 @@
fn f() {
var b = vec3<i32>();
var c = 0;
- let tint_symbol = c;
- let tint_symbol_1 = a(0);
- b[tint_symbol_1] = tint_symbol;
+ let tint_symbol = a(0);
+ b[tint_symbol] = c;
}
)";
@@ -3029,8 +3022,8 @@
fn f() {
var b = vec3<i32>();
var c = 0;
- let tint_symbol = a(0);
- b[c] = tint_symbol;
+ let tint_symbol = c;
+ b[tint_symbol] = a(0);
}
)";
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
index 2603bf4..d74a676 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
@@ -18,13 +18,13 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_3[4] = (int4[4])0;
- return tint_symbol_3;
+ const int4 tint_symbol_2[4] = (int4[4])0;
+ return tint_symbol_2;
}
S ret_struct_arr() {
- const S tint_symbol_4 = (S)0;
- return tint_symbol_4;
+ const S tint_symbol_3 = (S)0;
+ return tint_symbol_3;
}
void tint_symbol_store(uint offset, int4 value[4]) {
@@ -88,18 +88,17 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_5[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol_store(0u, tint_symbol_5);
+ const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol_store(0u, tint_symbol_4);
tint_symbol_store(0u, src_param);
- const int4 tint_symbol_1[4] = ret_arr();
- tint_symbol_store(0u, tint_symbol_1);
+ tint_symbol_store(0u, ret_arr());
const int4 src_let[4] = (int4[4])0;
tint_symbol_store(0u, src_let);
tint_symbol_store(0u, src_function);
tint_symbol_store(0u, src_private);
tint_symbol_store(0u, src_workgroup);
- const S tint_symbol_2 = ret_struct_arr();
- tint_symbol_store(0u, tint_symbol_2.arr);
+ const S tint_symbol_1 = ret_struct_arr();
+ tint_symbol_store(0u, tint_symbol_1.arr);
tint_symbol_store(0u, src_uniform_load(0u));
tint_symbol_store(0u, src_storage_load(0u));
int src_nested[4][3][2] = (int[4][3][2])0;
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
index 2603bf4..d74a676 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
@@ -18,13 +18,13 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_3[4] = (int4[4])0;
- return tint_symbol_3;
+ const int4 tint_symbol_2[4] = (int4[4])0;
+ return tint_symbol_2;
}
S ret_struct_arr() {
- const S tint_symbol_4 = (S)0;
- return tint_symbol_4;
+ const S tint_symbol_3 = (S)0;
+ return tint_symbol_3;
}
void tint_symbol_store(uint offset, int4 value[4]) {
@@ -88,18 +88,17 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_5[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol_store(0u, tint_symbol_5);
+ const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol_store(0u, tint_symbol_4);
tint_symbol_store(0u, src_param);
- const int4 tint_symbol_1[4] = ret_arr();
- tint_symbol_store(0u, tint_symbol_1);
+ tint_symbol_store(0u, ret_arr());
const int4 src_let[4] = (int4[4])0;
tint_symbol_store(0u, src_let);
tint_symbol_store(0u, src_function);
tint_symbol_store(0u, src_private);
tint_symbol_store(0u, src_workgroup);
- const S tint_symbol_2 = ret_struct_arr();
- tint_symbol_store(0u, tint_symbol_2.arr);
+ const S tint_symbol_1 = ret_struct_arr();
+ tint_symbol_store(0u, tint_symbol_1.arr);
tint_symbol_store(0u, src_uniform_load(0u));
tint_symbol_store(0u, src_storage_load(0u));
int src_nested[4][3][2] = (int[4][3][2])0;
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.glsl b/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
index 3d49ee2..41ab9c1 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
@@ -31,29 +31,28 @@
} dst_nested;
ivec4[4] ret_arr() {
- ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
- return tint_symbol_2;
+ ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ return tint_symbol_1;
}
S ret_struct_arr() {
- S tint_symbol_3 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
- return tint_symbol_3;
+ S tint_symbol_2 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
+ return tint_symbol_2;
}
void foo(ivec4 src_param[4]) {
ivec4 src_function[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
- ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
- dst.inner.arr = tint_symbol_4;
+ ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
+ dst.inner.arr = tint_symbol_3;
dst.inner.arr = src_param;
- ivec4 tint_symbol[4] = ret_arr();
- dst.inner.arr = tint_symbol;
+ dst.inner.arr = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
dst.inner.arr = src_let;
dst.inner.arr = src_function;
dst.inner.arr = src_private;
dst.inner.arr = src_workgroup;
- S tint_symbol_1 = ret_struct_arr();
- dst.inner.arr = tint_symbol_1.arr;
+ S tint_symbol = ret_struct_arr();
+ dst.inner.arr = tint_symbol.arr;
dst.inner.arr = src_uniform.inner.arr;
dst.inner.arr = src_storage.inner.arr;
int src_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index ee189f4..54070c6 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -23,33 +23,32 @@
};
tint_array<int4, 4> ret_arr() {
- tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
- return tint_symbol_2;
+ tint_array<int4, 4> const tint_symbol_1 = tint_array<int4, 4>{};
+ return tint_symbol_1;
}
S ret_struct_arr() {
- S const tint_symbol_3 = S{};
- return tint_symbol_3;
+ S const tint_symbol_2 = S{};
+ return tint_symbol_2;
}
-void foo(tint_array<int4, 4> src_param, device S* const tint_symbol_5, threadgroup tint_array<int4, 4>* const tint_symbol_7, const constant S* const tint_symbol_8, device S* const tint_symbol_9, device S_nested* const tint_symbol_10) {
- thread tint_array<int4, 4> tint_symbol_6 = {};
+void foo(tint_array<int4, 4> src_param, device S* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, device S_nested* const tint_symbol_9) {
+ thread tint_array<int4, 4> tint_symbol_5 = {};
tint_array<int4, 4> src_function = {};
- tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
- (*(tint_symbol_5)).arr = tint_symbol_4;
- (*(tint_symbol_5)).arr = src_param;
- tint_array<int4, 4> const tint_symbol = ret_arr();
- (*(tint_symbol_5)).arr = tint_symbol;
+ tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
+ (*(tint_symbol_4)).arr = tint_symbol_3;
+ (*(tint_symbol_4)).arr = src_param;
+ (*(tint_symbol_4)).arr = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
- (*(tint_symbol_5)).arr = src_let;
- (*(tint_symbol_5)).arr = src_function;
- (*(tint_symbol_5)).arr = tint_symbol_6;
- (*(tint_symbol_5)).arr = *(tint_symbol_7);
- S const tint_symbol_1 = ret_struct_arr();
- (*(tint_symbol_5)).arr = tint_symbol_1.arr;
- (*(tint_symbol_5)).arr = (*(tint_symbol_8)).arr;
- (*(tint_symbol_5)).arr = (*(tint_symbol_9)).arr;
+ (*(tint_symbol_4)).arr = src_let;
+ (*(tint_symbol_4)).arr = src_function;
+ (*(tint_symbol_4)).arr = tint_symbol_5;
+ (*(tint_symbol_4)).arr = *(tint_symbol_6);
+ S const tint_symbol = ret_struct_arr();
+ (*(tint_symbol_4)).arr = tint_symbol.arr;
+ (*(tint_symbol_4)).arr = (*(tint_symbol_7)).arr;
+ (*(tint_symbol_4)).arr = (*(tint_symbol_8)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
- (*(tint_symbol_10)).arr = src_nested;
+ (*(tint_symbol_9)).arr = src_nested;
}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm b/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
index 83ed0dd..c2fb753 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
@@ -114,9 +114,9 @@
OpStore %46 %53
%54 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
OpStore %54 %src_param
- %55 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
- %56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- OpStore %56 %55
+ %55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %56 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+ OpStore %55 %56
%57 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
OpStore %57 %8
%58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
diff --git a/test/tint/bug/tint/534.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/534.wgsl.expected.dxc.hlsl
index 11026c0..19231d7 100644
--- a/test/tint/bug/tint/534.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/534.wgsl.expected.dxc.hlsl
@@ -33,8 +33,8 @@
uint4 dstColorBits = uint4(dstColor);
{
for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
- const uint tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
- set_uint4(srcColorBits, i, tint_symbol_1);
+ const uint tint_symbol_1 = i;
+ set_uint4(srcColorBits, tint_symbol_1, ConvertToFp16FloatValue(srcColor[i]));
bool tint_tmp_1 = success;
if (tint_tmp_1) {
tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
diff --git a/test/tint/bug/tint/534.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/534.wgsl.expected.fxc.hlsl
index 11026c0..19231d7 100644
--- a/test/tint/bug/tint/534.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/534.wgsl.expected.fxc.hlsl
@@ -33,8 +33,8 @@
uint4 dstColorBits = uint4(dstColor);
{
for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
- const uint tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
- set_uint4(srcColorBits, i, tint_symbol_1);
+ const uint tint_symbol_1 = i;
+ set_uint4(srcColorBits, tint_symbol_1, ConvertToFp16FloatValue(srcColor[i]));
bool tint_tmp_1 = success;
if (tint_tmp_1) {
tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
diff --git a/test/tint/bug/tint/534.wgsl.expected.glsl b/test/tint/bug/tint/534.wgsl.expected.glsl
index 87a2177..8766dc1 100644
--- a/test/tint/bug/tint/534.wgsl.expected.glsl
+++ b/test/tint/bug/tint/534.wgsl.expected.glsl
@@ -35,8 +35,8 @@
uvec4 dstColorBits = uvec4(dstColor);
{
for(uint i = 0u; (i < uniforms.inner.channelCount); i = (i + 1u)) {
- uint tint_symbol_2 = ConvertToFp16FloatValue(srcColor[i]);
- srcColorBits[i] = tint_symbol_2;
+ uint tint_symbol_2 = i;
+ srcColorBits[tint_symbol_2] = ConvertToFp16FloatValue(srcColor[i]);
bool tint_tmp = success;
if (tint_tmp) {
tint_tmp = (srcColorBits[i] == dstColorBits[i]);
diff --git a/test/tint/bug/tint/534.wgsl.expected.msl b/test/tint/bug/tint/534.wgsl.expected.msl
index 2fa76ad..89afccd 100644
--- a/test/tint/bug/tint/534.wgsl.expected.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.msl
@@ -42,8 +42,8 @@
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);
for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
- uint const tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
- srcColorBits[i] = tint_symbol_1;
+ uint const tint_symbol_1 = i;
+ srcColorBits[tint_symbol_1] = ConvertToFp16FloatValue(srcColor[i]);
success = (success && (srcColorBits[i] == dstColorBits[i]));
}
uint outputIndex = ((GlobalInvocationID[1] * uint(size[0])) + GlobalInvocationID[0]);
diff --git a/test/tint/bug/tint/534.wgsl.expected.spvasm b/test/tint/bug/tint/534.wgsl.expected.spvasm
index 3f6f4a2..9b373fd 100644
--- a/test/tint/bug/tint/534.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/534.wgsl.expected.spvasm
@@ -170,13 +170,13 @@
%92 = OpLabel
OpBranch %82
%91 = OpLabel
- %94 = OpLoad %uint %i
- %96 = OpAccessChain %_ptr_Function_float %srcColor %94
- %97 = OpLoad %float %96
- %93 = OpFunctionCall %uint %ConvertToFp16FloatValue %97
- %98 = OpLoad %uint %i
- %99 = OpAccessChain %_ptr_Function_uint %srcColorBits %98
- OpStore %99 %93
+ %93 = OpLoad %uint %i
+ %94 = OpAccessChain %_ptr_Function_uint %srcColorBits %93
+ %96 = OpLoad %uint %i
+ %98 = OpAccessChain %_ptr_Function_float %srcColor %96
+ %99 = OpLoad %float %98
+ %95 = OpFunctionCall %uint %ConvertToFp16FloatValue %99
+ OpStore %94 %95
%100 = OpLoad %bool %success
OpSelectionMerge %101 None
OpBranchConditional %100 %102 %101
diff --git a/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
index f0d36c8..bfdda96 100644
--- a/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
@@ -47,7 +47,7 @@
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
-struct tint_symbol_3 {
+struct tint_symbol_5 {
uint3 local_id : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
uint3 global_id : SV_DispatchThreadID;
@@ -88,8 +88,9 @@
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRow + innerRow);
const uint inputCol = (tileColA + innerCol);
- const float tint_symbol = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
- mm_Asub[inputRow][inputCol] = tint_symbol;
+ const uint tint_symbol = inputRow;
+ const uint tint_symbol_1 = inputCol;
+ mm_Asub[tint_symbol][tint_symbol_1] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
}
}
}
@@ -100,8 +101,9 @@
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRowB + innerRow);
const uint inputCol = (tileCol + innerCol);
- const float tint_symbol_1 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
- mm_Bsub[innerCol][inputCol] = tint_symbol_1;
+ const uint tint_symbol_2 = innerCol;
+ const uint tint_symbol_3 = inputCol;
+ mm_Bsub[tint_symbol_2][tint_symbol_3] = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
}
}
}
@@ -143,7 +145,7 @@
}
[numthreads(16, 16, 1)]
-void main(tint_symbol_3 tint_symbol_2) {
- main_inner(tint_symbol_2.local_id, tint_symbol_2.global_id, tint_symbol_2.local_invocation_index);
+void main(tint_symbol_5 tint_symbol_4) {
+ main_inner(tint_symbol_4.local_id, tint_symbol_4.global_id, tint_symbol_4.local_invocation_index);
return;
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
index f0d36c8..bfdda96 100644
--- a/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
@@ -47,7 +47,7 @@
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
-struct tint_symbol_3 {
+struct tint_symbol_5 {
uint3 local_id : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
uint3 global_id : SV_DispatchThreadID;
@@ -88,8 +88,9 @@
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRow + innerRow);
const uint inputCol = (tileColA + innerCol);
- const float tint_symbol = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
- mm_Asub[inputRow][inputCol] = tint_symbol;
+ const uint tint_symbol = inputRow;
+ const uint tint_symbol_1 = inputCol;
+ mm_Asub[tint_symbol][tint_symbol_1] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
}
}
}
@@ -100,8 +101,9 @@
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRowB + innerRow);
const uint inputCol = (tileCol + innerCol);
- const float tint_symbol_1 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
- mm_Bsub[innerCol][inputCol] = tint_symbol_1;
+ const uint tint_symbol_2 = innerCol;
+ const uint tint_symbol_3 = inputCol;
+ mm_Bsub[tint_symbol_2][tint_symbol_3] = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
}
}
}
@@ -143,7 +145,7 @@
}
[numthreads(16, 16, 1)]
-void main(tint_symbol_3 tint_symbol_2) {
- main_inner(tint_symbol_2.local_id, tint_symbol_2.global_id, tint_symbol_2.local_invocation_index);
+void main(tint_symbol_5 tint_symbol_4) {
+ main_inner(tint_symbol_4.local_id, tint_symbol_4.global_id, tint_symbol_4.local_invocation_index);
return;
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.glsl b/test/tint/bug/tint/914.wgsl.expected.glsl
index cf748ec..f50c972 100644
--- a/test/tint/bug/tint/914.wgsl.expected.glsl
+++ b/test/tint/bug/tint/914.wgsl.expected.glsl
@@ -99,8 +99,9 @@
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
uint inputRow = (tileRow + innerRow);
uint inputCol = (tileColA + innerCol);
- float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
- mm_Asub[inputRow][inputCol] = tint_symbol_1;
+ uint tint_symbol_1 = inputRow;
+ uint tint_symbol_2 = inputCol;
+ mm_Asub[tint_symbol_1][tint_symbol_2] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
}
}
}
@@ -111,8 +112,9 @@
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint inputRow = (tileRowB + innerRow);
uint inputCol = (tileCol + innerCol);
- float tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
- mm_Bsub[innerCol][inputCol] = tint_symbol_2;
+ uint tint_symbol_3 = innerCol;
+ uint tint_symbol_4 = inputCol;
+ mm_Bsub[tint_symbol_3][tint_symbol_4] = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
}
}
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index acbc540..7dd3ccd 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -24,26 +24,26 @@
/* 0x0000 */ tint_array<float, 1> numbers;
};
-float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_3, const device Matrix* const tint_symbol_4) {
- if (((row < (*(tint_symbol_3)).dimAOuter) && (col < (*(tint_symbol_3)).dimInner))) {
- float const result = (*(tint_symbol_4)).numbers[((row * (*(tint_symbol_3)).dimInner) + col)];
+float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_5, const device Matrix* const tint_symbol_6) {
+ if (((row < (*(tint_symbol_5)).dimAOuter) && (col < (*(tint_symbol_5)).dimInner))) {
+ float const result = (*(tint_symbol_6)).numbers[((row * (*(tint_symbol_5)).dimInner) + col)];
return result;
}
return 0.0f;
}
-float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_5, const device Matrix* const tint_symbol_6) {
- if (((row < (*(tint_symbol_5)).dimInner) && (col < (*(tint_symbol_5)).dimBOuter))) {
- float const result = (*(tint_symbol_6)).numbers[((row * (*(tint_symbol_5)).dimBOuter) + col)];
+float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_7, const device Matrix* const tint_symbol_8) {
+ if (((row < (*(tint_symbol_7)).dimInner) && (col < (*(tint_symbol_7)).dimBOuter))) {
+ float const result = (*(tint_symbol_8)).numbers[((row * (*(tint_symbol_7)).dimBOuter) + col)];
return result;
}
return 0.0f;
}
-void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_7, device Matrix* const tint_symbol_8) {
- if (((row < (*(tint_symbol_7)).dimAOuter) && (col < (*(tint_symbol_7)).dimBOuter))) {
- uint const index = (col + (row * (*(tint_symbol_7)).dimBOuter));
- (*(tint_symbol_8)).numbers[index] = value;
+void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_9, device Matrix* const tint_symbol_10) {
+ if (((row < (*(tint_symbol_9)).dimAOuter) && (col < (*(tint_symbol_9)).dimBOuter))) {
+ uint const index = (col + (row * (*(tint_symbol_9)).dimBOuter));
+ (*(tint_symbol_10)).numbers[index] = value;
}
}
@@ -51,19 +51,19 @@
return (lhs / select(rhs, 1u, (rhs == 0u)));
}
-void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_9, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_10, const constant Uniforms* const tint_symbol_11, const device Matrix* const tint_symbol_12, const device Matrix* const tint_symbol_13, device Matrix* const tint_symbol_14) {
+void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_12, const constant Uniforms* const tint_symbol_13, const device Matrix* const tint_symbol_14, const device Matrix* const tint_symbol_15, device Matrix* const tint_symbol_16) {
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
uint const i = (idx / 64u);
uint const i_1 = (idx % 64u);
- (*(tint_symbol_9))[i][i_1] = 0.0f;
- (*(tint_symbol_10))[i][i_1] = 0.0f;
+ (*(tint_symbol_11))[i][i_1] = 0.0f;
+ (*(tint_symbol_12))[i][i_1] = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const tileRow = (local_id[1] * 4u);
uint const tileCol = (local_id[0] * 4u);
uint const globalRow = (global_id[1] * 4u);
uint const globalCol = (global_id[0] * 4u);
- uint const numTiles = (tint_div(((*(tint_symbol_11)).dimInner - 1u), 64u) + 1u);
+ uint const numTiles = (tint_div(((*(tint_symbol_13)).dimInner - 1u), 64u) + 1u);
tint_array<float, 16> acc = {};
float ACached = 0.0f;
tint_array<float, 4> BCached = {};
@@ -79,25 +79,27 @@
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRow + innerRow);
uint const inputCol = (tileColA + innerCol);
- float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_11, tint_symbol_12);
- (*(tint_symbol_9))[inputRow][inputCol] = tint_symbol_1;
+ uint const tint_symbol_1 = inputRow;
+ uint const tint_symbol_2 = inputCol;
+ (*(tint_symbol_11))[tint_symbol_1][tint_symbol_2] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_13, tint_symbol_14);
}
}
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRowB + innerRow);
uint const inputCol = (tileCol + innerCol);
- float const tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
- (*(tint_symbol_10))[innerCol][inputCol] = tint_symbol_2;
+ uint const tint_symbol_3 = innerCol;
+ uint const tint_symbol_4 = inputCol;
+ (*(tint_symbol_12))[tint_symbol_3][tint_symbol_4] = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_13, tint_symbol_15);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
for(uint k = 0u; (k < 64u); k = (k + 1u)) {
for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
- BCached[inner] = (*(tint_symbol_10))[k][(tileCol + inner)];
+ BCached[inner] = (*(tint_symbol_12))[k][(tileCol + inner)];
}
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
- ACached = (*(tint_symbol_9))[(tileRow + innerRow)][k];
+ ACached = (*(tint_symbol_11))[(tileRow + innerRow)][k];
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * 4u) + innerCol);
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
@@ -109,15 +111,15 @@
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * 4u) + innerCol);
- mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_11, tint_symbol_14);
+ mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_13, tint_symbol_16);
}
}
}
-kernel void tint_symbol(const constant Uniforms* tint_symbol_17 [[buffer(0)]], const device Matrix* tint_symbol_18 [[buffer(2)]], const device Matrix* tint_symbol_19 [[buffer(3)]], device Matrix* tint_symbol_20 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
- threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_15;
- threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_16;
- tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_15), &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
+kernel void tint_symbol(const constant Uniforms* tint_symbol_19 [[buffer(0)]], const device Matrix* tint_symbol_20 [[buffer(2)]], const device Matrix* tint_symbol_21 [[buffer(3)]], device Matrix* tint_symbol_22 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_17;
+ threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_18;
+ tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_17), &(tint_symbol_18), tint_symbol_19, tint_symbol_20, tint_symbol_21, tint_symbol_22);
return;
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.spvasm b/test/tint/bug/tint/914.wgsl.expected.spvasm
index 829ad7d..e45387e 100644
--- a/test/tint/bug/tint/914.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/914.wgsl.expected.spvasm
@@ -406,14 +406,14 @@
%228 = OpIAdd %uint %157 %227
%229 = OpLoad %uint %innerCol
%230 = OpIAdd %uint %194 %229
- %232 = OpLoad %uint %innerRow
- %233 = OpIAdd %uint %161 %232
- %234 = OpLoad %uint %t
- %235 = OpIMul %uint %234 %uint_64
- %236 = OpIAdd %uint %235 %230
- %231 = OpFunctionCall %float %mm_readA %233 %236
- %237 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %228 %230
- OpStore %237 %231
+ %231 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %228 %230
+ %233 = OpLoad %uint %innerRow
+ %234 = OpIAdd %uint %161 %233
+ %235 = OpLoad %uint %t
+ %236 = OpIMul %uint %235 %uint_64
+ %237 = OpIAdd %uint %236 %230
+ %232 = OpFunctionCall %float %mm_readA %234 %237
+ OpStore %231 %232
OpBranch %220
%220 = OpLabel
%238 = OpLoad %uint %innerCol
@@ -460,15 +460,15 @@
%263 = OpIAdd %uint %196 %262
%264 = OpLoad %uint %innerCol_0
%265 = OpIAdd %uint %159 %264
- %267 = OpLoad %uint %t
- %268 = OpIMul %uint %267 %uint_64
- %269 = OpIAdd %uint %268 %263
- %270 = OpLoad %uint %innerCol_0
- %271 = OpIAdd %uint %163 %270
- %266 = OpFunctionCall %float %mm_readB %269 %271
+ %266 = OpLoad %uint %innerCol_0
+ %267 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %266 %265
+ %269 = OpLoad %uint %t
+ %270 = OpIMul %uint %269 %uint_64
+ %271 = OpIAdd %uint %270 %263
%272 = OpLoad %uint %innerCol_0
- %273 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %272 %265
- OpStore %273 %266
+ %273 = OpIAdd %uint %163 %272
+ %268 = OpFunctionCall %float %mm_readB %271 %273
+ OpStore %267 %268
OpBranch %255
%255 = OpLabel
%274 = OpLoad %uint %innerCol_0
diff --git a/test/tint/bug/tint/980.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/980.wgsl.expected.dxc.hlsl
index 32fbf9d..2c3d0f4 100644
--- a/test/tint/bug/tint/980.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/980.wgsl.expected.dxc.hlsl
@@ -10,17 +10,16 @@
RWByteAddressBuffer io : register(u0, space0);
-struct tint_symbol_2 {
+struct tint_symbol_1 {
uint idx : SV_GroupIndex;
};
void main_inner(uint idx) {
- const float3 tint_symbol = Bad(io.Load(12u), asfloat(io.Load3(0u)));
- io.Store3(0u, asuint(tint_symbol));
+ io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
}
[numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
- main_inner(tint_symbol_1.idx);
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.idx);
return;
}
diff --git a/test/tint/bug/tint/980.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/980.wgsl.expected.fxc.hlsl
index 32fbf9d..2c3d0f4 100644
--- a/test/tint/bug/tint/980.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/980.wgsl.expected.fxc.hlsl
@@ -10,17 +10,16 @@
RWByteAddressBuffer io : register(u0, space0);
-struct tint_symbol_2 {
+struct tint_symbol_1 {
uint idx : SV_GroupIndex;
};
void main_inner(uint idx) {
- const float3 tint_symbol = Bad(io.Load(12u), asfloat(io.Load3(0u)));
- io.Store3(0u, asuint(tint_symbol));
+ io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
}
[numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
- main_inner(tint_symbol_1.idx);
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.idx);
return;
}
diff --git a/test/tint/bug/tint/980.wgsl.expected.glsl b/test/tint/bug/tint/980.wgsl.expected.glsl
index 8b59164..fc82f1aa 100644
--- a/test/tint/bug/tint/980.wgsl.expected.glsl
+++ b/test/tint/bug/tint/980.wgsl.expected.glsl
@@ -16,8 +16,7 @@
} io;
void tint_symbol(uint idx) {
- vec3 tint_symbol_1 = Bad(io.inner.i, io.inner.v);
- io.inner.v = tint_symbol_1;
+ io.inner.v = Bad(io.inner.i, io.inner.v);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/bug/tint/980.wgsl.expected.msl b/test/tint/bug/tint/980.wgsl.expected.msl
index cb9fabc..8d0305e 100644
--- a/test/tint/bug/tint/980.wgsl.expected.msl
+++ b/test/tint/bug/tint/980.wgsl.expected.msl
@@ -17,13 +17,12 @@
uint i;
};
-void tint_symbol_inner(uint idx, device S_tint_packed_vec3* const tint_symbol_2) {
- float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, float3((*(tint_symbol_2)).v));
- (*(tint_symbol_2)).v = packed_float3(tint_symbol_1);
+void tint_symbol_inner(uint idx, device S_tint_packed_vec3* const tint_symbol_1) {
+ (*(tint_symbol_1)).v = packed_float3(Bad((*(tint_symbol_1)).i, float3((*(tint_symbol_1)).v)));
}
-kernel void tint_symbol(device S_tint_packed_vec3* tint_symbol_3 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
- tint_symbol_inner(idx, tint_symbol_3);
+kernel void tint_symbol(device S_tint_packed_vec3* tint_symbol_2 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
+ tint_symbol_inner(idx, tint_symbol_2);
return;
}
diff --git a/test/tint/bug/tint/980.wgsl.expected.spvasm b/test/tint/bug/tint/980.wgsl.expected.spvasm
index 68c3650..a504c31 100644
--- a/test/tint/bug/tint/980.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/980.wgsl.expected.spvasm
@@ -45,9 +45,9 @@
%void = OpTypeVoid
%26 = OpTypeFunction %void %uint
%uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%uint_1 = OpConstant %uint 1
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
-%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%41 = OpTypeFunction %void
%Bad = OpFunction %v3float None %10
%index = OpFunctionParameter %uint
@@ -67,13 +67,13 @@
%main_inner = OpFunction %void None %26
%idx = OpFunctionParameter %uint
%30 = OpLabel
- %35 = OpAccessChain %_ptr_StorageBuffer_uint %io %uint_0 %uint_1
- %36 = OpLoad %uint %35
- %38 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0 %uint_0
- %39 = OpLoad %v3float %38
- %31 = OpFunctionCall %v3float %Bad %36 %39
- %40 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0 %uint_0
- OpStore %40 %31
+ %33 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0 %uint_0
+ %37 = OpAccessChain %_ptr_StorageBuffer_uint %io %uint_0 %uint_1
+ %38 = OpLoad %uint %37
+ %39 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0 %uint_0
+ %40 = OpLoad %v3float %39
+ %34 = OpFunctionCall %v3float %Bad %38 %40
+ OpStore %33 %34
OpReturn
OpFunctionEnd
%main = OpFunction %void None %41
diff --git a/test/tint/bug/tint/993.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/993.wgsl.expected.dxc.hlsl
index b08eb3f..7bf60d0 100644
--- a/test/tint/bug/tint/993.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/993.wgsl.expected.dxc.hlsl
@@ -20,7 +20,6 @@
[numthreads(1, 1, 1)]
void main() {
const int tint_symbol = runTest();
- const uint tint_symbol_1 = uint(tint_symbol);
- result.Store(0u, asuint(tint_symbol_1));
+ result.Store(0u, asuint(uint(tint_symbol)));
return;
}
diff --git a/test/tint/bug/tint/993.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/993.wgsl.expected.fxc.hlsl
index b08eb3f..7bf60d0 100644
--- a/test/tint/bug/tint/993.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/993.wgsl.expected.fxc.hlsl
@@ -20,7 +20,6 @@
[numthreads(1, 1, 1)]
void main() {
const int tint_symbol = runTest();
- const uint tint_symbol_1 = uint(tint_symbol);
- result.Store(0u, asuint(tint_symbol_1));
+ result.Store(0u, asuint(uint(tint_symbol)));
return;
}
diff --git a/test/tint/bug/tint/993.wgsl.expected.glsl b/test/tint/bug/tint/993.wgsl.expected.glsl
index 3c64f4e..3f237461 100644
--- a/test/tint/bug/tint/993.wgsl.expected.glsl
+++ b/test/tint/bug/tint/993.wgsl.expected.glsl
@@ -33,8 +33,7 @@
void tint_symbol() {
int tint_symbol_1 = runTest();
- uint tint_symbol_2 = uint(tint_symbol_1);
- result.inner.value = tint_symbol_2;
+ result.inner.value = uint(tint_symbol_1);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/bug/tint/993.wgsl.expected.msl b/test/tint/bug/tint/993.wgsl.expected.msl
index 7b82d95..b23a6ea 100644
--- a/test/tint/bug/tint/993.wgsl.expected.msl
+++ b/test/tint/bug/tint/993.wgsl.expected.msl
@@ -26,14 +26,13 @@
/* 0x0000 */ tint_array<atomic_int, 3> data;
};
-int runTest(device TestData* const tint_symbol_3, const constant Constants* const tint_symbol_4) {
- return atomic_load_explicit(&((*(tint_symbol_3)).data[(0u + uint((*(tint_symbol_4)).zero))]), memory_order_relaxed);
+int runTest(device TestData* const tint_symbol_2, const constant Constants* const tint_symbol_3) {
+ return atomic_load_explicit(&((*(tint_symbol_2)).data[(0u + uint((*(tint_symbol_3)).zero))]), memory_order_relaxed);
}
-kernel void tint_symbol(device TestData* tint_symbol_5 [[buffer(2)]], const constant Constants* tint_symbol_6 [[buffer(0)]], device Result* tint_symbol_7 [[buffer(1)]]) {
- int const tint_symbol_1 = runTest(tint_symbol_5, tint_symbol_6);
- uint const tint_symbol_2 = uint(tint_symbol_1);
- (*(tint_symbol_7)).value = tint_symbol_2;
+kernel void tint_symbol(device TestData* tint_symbol_4 [[buffer(2)]], const constant Constants* tint_symbol_5 [[buffer(0)]], device Result* tint_symbol_6 [[buffer(1)]]) {
+ int const tint_symbol_1 = runTest(tint_symbol_4, tint_symbol_5);
+ (*(tint_symbol_6)).value = uint(tint_symbol_1);
return;
}
diff --git a/test/tint/bug/tint/993.wgsl.expected.spvasm b/test/tint/bug/tint/993.wgsl.expected.spvasm
index 9c13129..b95c928 100644
--- a/test/tint/bug/tint/993.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/993.wgsl.expected.spvasm
@@ -78,8 +78,8 @@
%main = OpFunction %void None %32
%35 = OpLabel
%36 = OpFunctionCall %int %runTest
- %37 = OpBitcast %uint %36
- %39 = OpAccessChain %_ptr_StorageBuffer_uint %result %uint_0 %uint_0
- OpStore %39 %37
+ %38 = OpAccessChain %_ptr_StorageBuffer_uint %result %uint_0 %uint_0
+ %39 = OpBitcast %uint %36
+ OpStore %38 %39
OpReturn
OpFunctionEnd
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.dxc.hlsl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.dxc.hlsl
index 45f02cd..d073ee1 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.dxc.hlsl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.dxc.hlsl
@@ -10,15 +10,15 @@
typedef strided_arr mat2x2_stride_16_to_arr_ret[2];
mat2x2_stride_16_to_arr_ret mat2x2_stride_16_to_arr(float2x2 m) {
- const strided_arr tint_symbol_1 = {m[0u]};
- const strided_arr tint_symbol_2 = {m[1u]};
- const strided_arr tint_symbol_3[2] = {tint_symbol_1, tint_symbol_2};
- return tint_symbol_3;
+ const strided_arr tint_symbol = {m[0u]};
+ const strided_arr tint_symbol_1 = {m[1u]};
+ const strided_arr tint_symbol_2[2] = {tint_symbol, tint_symbol_1};
+ return tint_symbol_2;
}
strided_arr ssbo_load_1(uint offset) {
- const strided_arr tint_symbol_4 = {asfloat(ssbo.Load2((offset + 0u)))};
- return tint_symbol_4;
+ const strided_arr tint_symbol_3 = {asfloat(ssbo.Load2((offset + 0u)))};
+ return tint_symbol_3;
}
typedef strided_arr ssbo_load_ret[2];
@@ -47,8 +47,7 @@
void f_1() {
const float2x2 x_15 = arr_to_mat2x2_stride_16(ssbo_load(0u));
- const strided_arr tint_symbol[2] = mat2x2_stride_16_to_arr(x_15);
- ssbo_store(0u, tint_symbol);
+ ssbo_store(0u, mat2x2_stride_16_to_arr(x_15));
return;
}
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.fxc.hlsl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.fxc.hlsl
index 45f02cd..d073ee1 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.fxc.hlsl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.fxc.hlsl
@@ -10,15 +10,15 @@
typedef strided_arr mat2x2_stride_16_to_arr_ret[2];
mat2x2_stride_16_to_arr_ret mat2x2_stride_16_to_arr(float2x2 m) {
- const strided_arr tint_symbol_1 = {m[0u]};
- const strided_arr tint_symbol_2 = {m[1u]};
- const strided_arr tint_symbol_3[2] = {tint_symbol_1, tint_symbol_2};
- return tint_symbol_3;
+ const strided_arr tint_symbol = {m[0u]};
+ const strided_arr tint_symbol_1 = {m[1u]};
+ const strided_arr tint_symbol_2[2] = {tint_symbol, tint_symbol_1};
+ return tint_symbol_2;
}
strided_arr ssbo_load_1(uint offset) {
- const strided_arr tint_symbol_4 = {asfloat(ssbo.Load2((offset + 0u)))};
- return tint_symbol_4;
+ const strided_arr tint_symbol_3 = {asfloat(ssbo.Load2((offset + 0u)))};
+ return tint_symbol_3;
}
typedef strided_arr ssbo_load_ret[2];
@@ -47,8 +47,7 @@
void f_1() {
const float2x2 x_15 = arr_to_mat2x2_stride_16(ssbo_load(0u));
- const strided_arr tint_symbol[2] = mat2x2_stride_16_to_arr(x_15);
- ssbo_store(0u, tint_symbol);
+ ssbo_store(0u, mat2x2_stride_16_to_arr(x_15));
return;
}
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl
index 5be0805..2299089 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl
@@ -32,9 +32,8 @@
void tint_symbol() {
InnerS v = InnerS(0);
OuterS s = OuterS(S1[8](S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0)))));
- InnerS tint_symbol_1 = v;
- uint tint_symbol_2 = getNextIndex();
- s.a1[tint_symbol_2].a2[uniforms.inner.j] = tint_symbol_1;
+ uint tint_symbol_1 = getNextIndex();
+ s.a1[tint_symbol_1].a2[uniforms.inner.j] = v;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl
index 7b606ba..4cd6b66 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl
@@ -32,17 +32,16 @@
};
uint getNextIndex() {
- thread uint tint_symbol_3 = 0u;
- tint_symbol_3 = (tint_symbol_3 + 1u);
- return tint_symbol_3;
+ thread uint tint_symbol_2 = 0u;
+ tint_symbol_2 = (tint_symbol_2 + 1u);
+ return tint_symbol_2;
}
-kernel void tint_symbol(const constant Uniforms* tint_symbol_4 [[buffer(0)]]) {
+kernel void tint_symbol(const constant Uniforms* tint_symbol_3 [[buffer(0)]]) {
InnerS v = {};
OuterS s = {};
- InnerS const tint_symbol_1 = v;
- uint const tint_symbol_2 = getNextIndex();
- s.a1[tint_symbol_2].a2[(*(tint_symbol_4)).j] = tint_symbol_1;
+ uint const tint_symbol_1 = getNextIndex();
+ s.a1[tint_symbol_1].a2[(*(tint_symbol_3)).j] = v;
return;
}
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm
index a9aa39b..2047796 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm
@@ -73,11 +73,11 @@
%19 = OpLabel
%v = OpVariable %_ptr_Function_InnerS Function %24
%s = OpVariable %_ptr_Function_OuterS Function %32
- %33 = OpLoad %InnerS %v
- %34 = OpFunctionCall %uint %getNextIndex
- %37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %38 = OpLoad %uint %37
- %39 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %34 %uint_0 %38
- OpStore %39 %33
+ %33 = OpFunctionCall %uint %getNextIndex
+ %36 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %37 = OpLoad %uint %36
+ %38 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %37
+ %39 = OpLoad %InnerS %v
+ OpStore %38 %39
OpReturn
OpFunctionEnd
diff --git a/webgpu-cts/expectations.txt b/webgpu-cts/expectations.txt
index 54b1a59..6c5b7c9 100644
--- a/webgpu-cts/expectations.txt
+++ b/webgpu-cts/expectations.txt
@@ -69,6 +69,11 @@
# Last rolled: 2023-03-03 12:32:02AM
################################################################################
+# Evaluation order for assignments is changing
+################################################################################
+crbug.com/tint/1867 webgpu:shader,execution,evaluation_order:assignment:* [ Failure ]
+
+################################################################################
# copyToTexture failures on Linux
# Skipped instead of just Crash because of the number of failures
# KEEP
