encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Ben Clayton 2023-03-08 21:48:45 +00:00 committed by Dawn LUCI CQ
parent 43c5efa7e8
commit da353b4b39
35 changed files with 432 additions and 293 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

80
src/tint/resolver/uniformity.cc
View File

@ -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 @@ 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<Node*, 4> edges;
@ -547,14 +551,15 @@ class UniformityGraph {
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 @@ 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,11 +1374,23 @@ 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<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,
LValue ProcessLValueExpression(Node* cf,
const ast::Expression* expr,
bool is_partial_reference = false) {
return Switch(
@ -1378,35 +1399,37 @@ class UniformityGraph {
[&](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 @@ 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<Node*, Node*>(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<Node*, Node*>(nullptr, nullptr);
return LValue{};
});
}

146
src/tint/resolver/uniformity_test.cc
View File

@ -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<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

57
src/tint/transform/promote_side_effects_to_decl.cc
View File

@ -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<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 @@ 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<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 @@ class DecomposeSideEffects::DecomposeState : public StateBase {
!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);
},

43
src/tint/transform/promote_side_effects_to_decl_test.cc
View File

@ -2830,9 +2830,8 @@ fn a(i : i32) -> i32 {
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 a(i : i32) -> i32 {
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 a(i : i32) -> i32 {
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 @@ fn f() {
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 a(i : i32) -> i32 {
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 a(i : i32) -> i32 {
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 a(i : i32) -> i32 {
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);
}
)";

19
test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
View File

@ -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;

19
test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
View File

@ -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;

19
test/tint/array/assign_to_storage_var.wgsl.expected.glsl
View File

@ -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)));

39
test/tint/array/assign_to_storage_var.wgsl.expected.msl
View File

@ -23,33 +23,32 @@ struct S_nested {
};
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;
}

6
test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
View File

@ -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

4
test/tint/bug/tint/534.wgsl.expected.dxc.hlsl
View File

@ -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]);

4
test/tint/bug/tint/534.wgsl.expected.fxc.hlsl
View File

@ -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]);

4
test/tint/bug/tint/534.wgsl.expected.glsl
View File

@ -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]);

4
test/tint/bug/tint/534.wgsl.expected.msl
View File

@ -42,8 +42,8 @@ void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample
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]);

14
test/tint/bug/tint/534.wgsl.expected.spvasm
View File

@ -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

16
test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
View File

@ -47,7 +47,7 @@ uint tint_div(uint lhs, uint rhs) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
}
[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;
}

16
test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
View File

@ -47,7 +47,7 @@ uint tint_div(uint lhs, uint rhs) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
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 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
}
[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;
}

10
test/tint/bug/tint/914.wgsl.expected.glsl
View File

@ -99,8 +99,9 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
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 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
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));
}
}
}

52
test/tint/bug/tint/914.wgsl.expected.msl
View File

@ -24,26 +24,26 @@ struct Matrix {
/* 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 @@ 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<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 @@ 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<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;
}

32
test/tint/bug/tint/914.wgsl.expected.spvasm
View File

@ -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

9
test/tint/bug/tint/980.wgsl.expected.dxc.hlsl
View File

@ -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;
}

9
test/tint/bug/tint/980.wgsl.expected.fxc.hlsl
View File

@ -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;
}

3
test/tint/bug/tint/980.wgsl.expected.glsl
View File

@ -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;

9
test/tint/bug/tint/980.wgsl.expected.msl
View File

@ -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;
}

16
test/tint/bug/tint/980.wgsl.expected.spvasm
View File

@ -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

3
test/tint/bug/tint/993.wgsl.expected.dxc.hlsl
View File

@ -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;
}

3
test/tint/bug/tint/993.wgsl.expected.fxc.hlsl
View File

@ -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;
}

3
test/tint/bug/tint/993.wgsl.expected.glsl
View File

@ -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;

11
test/tint/bug/tint/993.wgsl.expected.msl
View File

@ -26,14 +26,13 @@ struct TestData {
/* 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;
}

6
test/tint/bug/tint/993.wgsl.expected.spvasm
View File

@ -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

15
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.dxc.hlsl
View File

@ -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;
}

15
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.fxc.hlsl
View File

@ -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;
}

5
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl
View File

@ -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;

13
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl
View File

@ -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;
}

12
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm
View File

@ -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

5
webgpu-cts/expectations.txt
View File

@ -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