diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc index bc4ac1d281..673c1c0a0d 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 +#endif + namespace tint::resolver { namespace { @@ -123,7 +127,7 @@ struct Node { 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 edges; @@ -547,14 +551,15 @@ class UniformityGraph { stmt, [&](const ast::AssignmentStatement* a) { - auto [cf1, v1] = ProcessExpression(cf, a->rhs); if (a->lhs->Is()) { - 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 @@ class UniformityGraph { }, [&](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 @@ class UniformityGraph { 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 @@ class UniformityGraph { 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 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 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(); if (sem->Variable()->Is()) { - 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()) { // 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(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(cf2, l1); + return LValue{cf2, l1, apply}; }, [&](const ast::MemberAccessorExpression* m) { @@ -1419,17 +1442,18 @@ class UniformityGraph { // 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(cf, deref); + return LValue{cf, deref, apply}; } return ProcessLValueExpression(cf, u->expr, is_partial_reference); }, @@ -1437,7 +1461,7 @@ class UniformityGraph { [&](Default) { TINT_ICE(Resolver, diagnostics_) << "unknown lvalue expression type: " << std::string(expr->TypeInfo().name); - return std::pair(nullptr, nullptr); + return LValue{}; }); } diff --git a/src/tint/resolver/uniformity_test.cc b/src/tint/resolver/uniformity_test.cc index 51608bdb78..a55d259d8a 100644 --- a/src/tint/resolver/uniformity_test.cc +++ b/src/tint/resolver/uniformity_test.cc @@ -8503,5 +8503,151 @@ test:11:7 note: reading from read_write storage buffer 'non_uniform' may result )"); } +TEST_F(UniformityAnalysisTest, AssignmentEval_LHS_Then_RHS_Pass) { + std::string src = R"( +@group(0) @binding(0) var non_uniform : i32; + +fn b(p : ptr) -> i32 { + *p = non_uniform; + return 0; +} + +fn a(p : ptr) -> i32 { + if (*p == 0) { + workgroupBarrier(); + } + return 0; +} + +fn foo() { + var i = 0; + var arr : array; + 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 non_uniform : i32; + +fn a(p : ptr) -> i32 { + *p = non_uniform; + return 0; +} + +fn b(p : ptr) -> i32 { + if (*p == 0) { + workgroupBarrier(); + } + return 0; +} + +fn foo() { + var i = 0; + var arr : array; + 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 non_uniform : i32; + +fn b(p : ptr) -> i32 { + *p = non_uniform; + return 0; +} + +fn a(p : ptr) -> i32 { + if (*p == 0) { + workgroupBarrier(); + } + return 0; +} + +fn foo() { + var i = 0; + var arr : array; + 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 non_uniform : i32; + +fn a(p : ptr) -> i32 { + *p = non_uniform; + return 0; +} + +fn b(p : ptr) -> i32 { + if (*p == 0) { + workgroupBarrier(); + } + return 0; +} + +fn foo() { + var i = 0; + var arr : array; + 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 65ec7316cc..9db7e388e1 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 @@ class DecomposeSideEffects::CollectHoistsState : public StateBase { }); } - // 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 @@ class DecomposeSideEffects::CollectHoistsState : public StateBase { 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 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()) { - 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 @@ class DecomposeSideEffects::CollectHoistsState : public StateBase { } 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 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 @@ class DecomposeSideEffects::DecomposeState : public StateBase { !sem.GetVal(s->rhs)->HasSideEffects()) { return nullptr; } - // rhs before lhs + // lhs before rhs tint::utils::Vector 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 6ed4b27c87..636a8099fa 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 a(i : i32) -> i32 { fn f() { var b = array(); - 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 a(i : i32) -> i32 { fn f() { var b = array, 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 a(i : i32) -> i32 { fn f() { var b = array, 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 @@ fn f() { var b = array(); var d = array, 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 a(i : i32) -> i32 { fn f() { var b = vec3(); - 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 a(i : i32) -> i32 { fn f() { var b = vec3(); 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 a(i : i32) -> i32 { fn f() { var b = vec3(); 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 2603bf4e1f..d74a6769b7 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 @@ RWByteAddressBuffer dst_nested : register(u3, space0); 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 dst_nested_store(uint offset, int value[4][3][2]) { 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 2603bf4e1f..d74a6769b7 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 @@ RWByteAddressBuffer dst_nested : register(u3, space0); 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 dst_nested_store(uint offset, int value[4][3][2]) { 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 3d49ee2056..41ab9c1931 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 @@ layout(binding = 3, std430) buffer dst_nested_block_ssbo { } 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 ee189f41d7..54070c628e 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 @@ struct S_nested { }; tint_array ret_arr() { - tint_array const tint_symbol_2 = tint_array{}; - return tint_symbol_2; + tint_array const tint_symbol_1 = tint_array{}; + 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 src_param, device S* const tint_symbol_5, threadgroup tint_array* 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 tint_symbol_6 = {}; +void foo(tint_array src_param, device S* const tint_symbol_4, threadgroup tint_array* 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 tint_symbol_5 = {}; tint_array src_function = {}; - tint_array const tint_symbol_4 = tint_array{int4(1), int4(2), int4(3), int4(3)}; - (*(tint_symbol_5)).arr = tint_symbol_4; - (*(tint_symbol_5)).arr = src_param; - tint_array const tint_symbol = ret_arr(); - (*(tint_symbol_5)).arr = tint_symbol; + tint_array const tint_symbol_3 = tint_array{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 const src_let = tint_array{}; - (*(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, 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 83ed0dd821..c2fb753446 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 11026c0f57..19231d71f2 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 @@ void main_inner(uint3 GlobalInvocationID) { 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 11026c0f57..19231d71f2 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 @@ void main_inner(uint3 GlobalInvocationID) { 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 87a2177a65..8766dc1b7f 100644 --- a/test/tint/bug/tint/534.wgsl.expected.glsl +++ b/test/tint/bug/tint/534.wgsl.expected.glsl @@ -35,8 +35,8 @@ void tint_symbol_1(uvec3 GlobalInvocationID) { 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 2fa76ad477..89afccdfa3 100644 --- a/test/tint/bug/tint/534.wgsl.expected.msl +++ b/test/tint/bug/tint/534.wgsl.expected.msl @@ -42,8 +42,8 @@ void tint_symbol_inner(uint3 GlobalInvocationID, texture2d 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 @@ uint tint_div(uint lhs, uint rhs) { return (lhs / select(rhs, 1u, (rhs == 0u))); } -void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array, 64>* const tint_symbol_9, threadgroup tint_array, 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, 64>* const tint_symbol_11, threadgroup tint_array, 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 acc = {}; float ACached = 0.0f; tint_array BCached = {}; @@ -79,25 +79,27 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in 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 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in 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, 64> tint_symbol_15; - threadgroup tint_array, 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, 64> tint_symbol_17; + threadgroup tint_array, 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 829ad7d69f..e45387e03e 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 32fbf9d156..2c3d0f4402 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 @@ float3 Bad(uint index, float3 rd) { 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 32fbf9d156..2c3d0f4402 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 @@ float3 Bad(uint index, float3 rd) { 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 8b59164ab7..fc82f1aa43 100644 --- a/test/tint/bug/tint/980.wgsl.expected.glsl +++ b/test/tint/bug/tint/980.wgsl.expected.glsl @@ -16,8 +16,7 @@ layout(binding = 0, std430) buffer io_block_ssbo { } 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 cb9fabca5a..8d0305e395 100644 --- a/test/tint/bug/tint/980.wgsl.expected.msl +++ b/test/tint/bug/tint/980.wgsl.expected.msl @@ -17,13 +17,12 @@ struct S { 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 68c365025b..a504c318f9 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 b08eb3fd1a..7bf60d00cc 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 @@ int runTest() { [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 b08eb3fd1a..7bf60d00cc 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 @@ int runTest() { [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 3c64f4e965..3f23746182 100644 --- a/test/tint/bug/tint/993.wgsl.expected.glsl +++ b/test/tint/bug/tint/993.wgsl.expected.glsl @@ -33,8 +33,7 @@ int runTest() { 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 7b82d9534f..b23a6ea7e7 100644 --- a/test/tint/bug/tint/993.wgsl.expected.msl +++ b/test/tint/bug/tint/993.wgsl.expected.msl @@ -26,14 +26,13 @@ struct TestData { /* 0x0000 */ tint_array 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 9c131290cb..b95c928a77 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 45f02cd634..d073ee136a 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 @@ float2x2 arr_to_mat2x2_stride_16(strided_arr arr[2]) { 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 ssbo_store(uint offset, strided_arr value[2]) { 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 45f02cd634..d073ee136a 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 @@ float2x2 arr_to_mat2x2_stride_16(strided_arr arr[2]) { 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 ssbo_store(uint offset, strided_arr value[2]) { 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 5be0805cb3..2299089ed3 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 @@ layout(binding = 4, std140) uniform uniforms_block_ubo { 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 7b606ba7e6..4cd6b6630c 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 @@ struct OuterS { }; 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 a9aa39b1dd..2047796bbf 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 54b1a59a8f..6c5b7c9c6e 100644 --- a/webgpu-cts/expectations.txt +++ b/webgpu-cts/expectations.txt @@ -68,6 +68,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