Import Tint changes from Dawn
Changes:
- fbabce4ec77f087877e12e5ee9b3e0dea80ec08d [tint][ir] Validate no calls to entry points by James Price <jrprice@google.com>
- ef011009ed1479b5dc554f26a1e7da31cd957f79 [tint][wgsl][reader] Expand expression Source spans by Ben Clayton <bclayton@google.com>
- 842ab015bacf7d32fa978db60b0ff9c4075d79fe [tint][uniformity] Fix mutable pointer parameters by James Price <jrprice@google.com>
- af25bec4a0ec2804567e18fe1feaf817f18c38f4 [tint][ast] Fix dynamic indices in DirectVariableAccess by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: fbabce4ec77f087877e12e5ee9b3e0dea80ec08d
Change-Id: If0f9d64c7f1f3abbf37d82182caf3414609af33c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/169800
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access.cc b/src/tint/lang/core/ir/transform/direct_variable_access.cc
index 12b078b..e62d52b 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access.cc
@@ -380,7 +380,7 @@
return variant_fn;
});
- // Repoint the target of the call to the variant.
+ // Re-point the target of the call to the variant.
call->SetTarget(new_target);
});
}
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access_test.cc b/src/tint/lang/core/ir/transform/direct_variable_access_test.cc
index ed37b51..e6244da 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access_test.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access_test.cc
@@ -1158,6 +1158,128 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_DirectVariableAccessTest_UniformAS, CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ Var* input = nullptr;
+ b.Append(b.ir.root_block,
+ [&] { //
+ input = b.Var("U", ty.ptr<uniform>(T));
+ input->SetBindingPoint(0, 0);
+ });
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<uniform>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(),
+ [&] { b.Return(f2, b.Load(b.Access<ptr<uniform, vec4<i32>>>(p, 3_u))); });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<uniform>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<uniform>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<uniform>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<uniform>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ b.Call(f0, input);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%b1 = block { # root
+ %U:ptr<uniform, array<array<array<vec4<i32>, 5>, 5>, 5>, read> = var @binding_point(0, 0)
+}
+
+%f2 = func(%p:ptr<uniform, array<vec4<i32>, 5>, read>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:ptr<uniform, vec4<i32>, read> = access %p, 3u
+ %5:vec4<i32> = load %4
+ ret %5
+ }
+}
+%f1 = func(%p_1:ptr<uniform, array<array<vec4<i32>, 5>, 5>, read>):vec4<i32> -> %b3 { # %p_1: 'p'
+ %b3 = block {
+ %8:ptr<uniform, array<vec4<i32>, 5>, read> = access %p_1, 2u
+ %9:vec4<i32> = call %f2, %8
+ ret %9
+ }
+}
+%f0 = func(%p_2:ptr<uniform, array<array<array<vec4<i32>, 5>, 5>, 5>, read>):vec4<i32> -> %b4 { # %p_2: 'p'
+ %b4 = block {
+ %12:ptr<uniform, array<array<vec4<i32>, 5>, 5>, read> = access %p_2, 1u
+ %13:vec4<i32> = call %f1, %12
+ ret %13
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %15:vec4<i32> = call %f0, %U
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %U:ptr<uniform, array<array<array<vec4<i32>, 5>, 5>, 5>, read> = var @binding_point(0, 0)
+}
+
+%f2_U_X_X = func(%p_indices:array<u32, 2>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<uniform, array<vec4<i32>, 5>, read> = access %U, %4, %5
+ %7:ptr<uniform, vec4<i32>, read> = access %6, 3u
+ %8:vec4<i32> = load %7
+ ret %8
+ }
+}
+%f1_U_X = func(%p_indices_1:array<u32, 1>):vec4<i32> -> %b3 { # %p_indices_1: 'p_indices'
+ %b3 = block {
+ %11:u32 = access %p_indices_1, 0u
+ %12:array<u32, 2> = construct %11, 2u
+ %13:vec4<i32> = call %f2_U_X_X, %12
+ ret %13
+ }
+}
+%f0_U = func():vec4<i32> -> %b4 {
+ %b4 = block {
+ %15:array<u32, 1> = construct 1u
+ %16:vec4<i32> = call %f1_U_X, %15
+ ret %16
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %18:vec4<i32> = call %f0_U
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace uniform_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -1727,6 +1849,128 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_DirectVariableAccessTest_StorageAS, CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ Var* input = nullptr;
+ b.Append(b.ir.root_block,
+ [&] { //
+ input = b.Var("U", ty.ptr<storage>(T));
+ input->SetBindingPoint(0, 0);
+ });
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<storage>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(),
+ [&] { b.Return(f2, b.Load(b.Access<ptr<storage, vec4<i32>>>(p, 3_u))); });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<storage>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<storage>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<storage>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<storage>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ b.Call(f0, input);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%b1 = block { # root
+ %U:ptr<storage, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+}
+
+%f2 = func(%p:ptr<storage, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:ptr<storage, vec4<i32>, read_write> = access %p, 3u
+ %5:vec4<i32> = load %4
+ ret %5
+ }
+}
+%f1 = func(%p_1:ptr<storage, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_1: 'p'
+ %b3 = block {
+ %8:ptr<storage, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %9:vec4<i32> = call %f2, %8
+ ret %9
+ }
+}
+%f0 = func(%p_2:ptr<storage, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b4 { # %p_2: 'p'
+ %b4 = block {
+ %12:ptr<storage, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %13:vec4<i32> = call %f1, %12
+ ret %13
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %15:vec4<i32> = call %f0, %U
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %U:ptr<storage, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+}
+
+%f2_U_X_X = func(%p_indices:array<u32, 2>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<storage, array<vec4<i32>, 5>, read_write> = access %U, %4, %5
+ %7:ptr<storage, vec4<i32>, read_write> = access %6, 3u
+ %8:vec4<i32> = load %7
+ ret %8
+ }
+}
+%f1_U_X = func(%p_indices_1:array<u32, 1>):vec4<i32> -> %b3 { # %p_indices_1: 'p_indices'
+ %b3 = block {
+ %11:u32 = access %p_indices_1, 0u
+ %12:array<u32, 2> = construct %11, 2u
+ %13:vec4<i32> = call %f2_U_X_X, %12
+ ret %13
+ }
+}
+%f0_U = func():vec4<i32> -> %b4 {
+ %b4 = block {
+ %15:array<u32, 1> = construct 1u
+ %16:vec4<i32> = call %f1_U_X, %15
+ ret %16
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %18:vec4<i32> = call %f0_U
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace storage_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -2194,6 +2438,128 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_DirectVariableAccessTest_WorkgroupAS, CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ Var* input = nullptr;
+ b.Append(b.ir.root_block,
+ [&] { //
+ input = b.Var("U", ty.ptr<workgroup>(T));
+ input->SetBindingPoint(0, 0);
+ });
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<workgroup>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(),
+ [&] { b.Return(f2, b.Load(b.Access<ptr<workgroup, vec4<i32>>>(p, 3_u))); });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<workgroup>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<workgroup>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<workgroup>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<workgroup>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ b.Call(f0, input);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%b1 = block { # root
+ %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+}
+
+%f2 = func(%p:ptr<workgroup, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:ptr<workgroup, vec4<i32>, read_write> = access %p, 3u
+ %5:vec4<i32> = load %4
+ ret %5
+ }
+}
+%f1 = func(%p_1:ptr<workgroup, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_1: 'p'
+ %b3 = block {
+ %8:ptr<workgroup, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %9:vec4<i32> = call %f2, %8
+ ret %9
+ }
+}
+%f0 = func(%p_2:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b4 { # %p_2: 'p'
+ %b4 = block {
+ %12:ptr<workgroup, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %13:vec4<i32> = call %f1, %12
+ ret %13
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %15:vec4<i32> = call %f0, %U
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+}
+
+%f2_U_X_X = func(%p_indices:array<u32, 2>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<workgroup, array<vec4<i32>, 5>, read_write> = access %U, %4, %5
+ %7:ptr<workgroup, vec4<i32>, read_write> = access %6, 3u
+ %8:vec4<i32> = load %7
+ ret %8
+ }
+}
+%f1_U_X = func(%p_indices_1:array<u32, 1>):vec4<i32> -> %b3 { # %p_indices_1: 'p_indices'
+ %b3 = block {
+ %11:u32 = access %p_indices_1, 0u
+ %12:array<u32, 2> = construct %11, 2u
+ %13:vec4<i32> = call %f2_U_X_X, %12
+ ret %13
+ }
+}
+%f0_U = func():vec4<i32> -> %b4 {
+ %b4 = block {
+ %15:array<u32, 1> = construct 1u
+ %16:vec4<i32> = call %f1_U_X, %15
+ ret %16
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %18:vec4<i32> = call %f0_U
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace workgroup_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -3323,6 +3689,214 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_DirectVariableAccessTest_PrivateAS, Enabled_CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ Var* P = nullptr;
+ b.Append(b.ir.root_block,
+ [&] { //
+ P = b.Var("P", ty.ptr<private_>(T));
+ });
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ b.Return(f2, b.Load(b.Access<ptr<private_, vec4<i32>, read_write>>(p, 3_u)));
+ });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<private_>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<private_>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ b.Call(f0, P);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%b1 = block { # root
+ %P:ptr<private, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+}
+
+%f2 = func(%p:ptr<private, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:ptr<private, vec4<i32>, read_write> = access %p, 3u
+ %5:vec4<i32> = load %4
+ ret %5
+ }
+}
+%f1 = func(%p_1:ptr<private, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_1: 'p'
+ %b3 = block {
+ %8:ptr<private, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %9:vec4<i32> = call %f2, %8
+ ret %9
+ }
+}
+%f0 = func(%p_2:ptr<private, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b4 { # %p_2: 'p'
+ %b4 = block {
+ %12:ptr<private, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %13:vec4<i32> = call %f1, %12
+ ret %13
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %15:vec4<i32> = call %f0, %P
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %P:ptr<private, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+}
+
+%f2_P_X_X = func(%p_indices:array<u32, 2>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<private, array<vec4<i32>, 5>, read_write> = access %P, %4, %5
+ %7:ptr<private, vec4<i32>, read_write> = access %6, 3u
+ %8:vec4<i32> = load %7
+ ret %8
+ }
+}
+%f1_P_X = func(%p_indices_1:array<u32, 1>):vec4<i32> -> %b3 { # %p_indices_1: 'p_indices'
+ %b3 = block {
+ %11:u32 = access %p_indices_1, 0u
+ %12:array<u32, 2> = construct %11, 2u
+ %13:vec4<i32> = call %f2_P_X_X, %12
+ ret %13
+ }
+}
+%f0_P = func():vec4<i32> -> %b4 {
+ %b4 = block {
+ %15:array<u32, 1> = construct 1u
+ %16:vec4<i32> = call %f1_P_X, %15
+ ret %16
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %18:vec4<i32> = call %f0_P
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, kTransformPrivate);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_DirectVariableAccessTest_PrivateAS, Disabled_CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ Var* P = nullptr;
+ b.Append(b.ir.root_block,
+ [&] { //
+ P = b.Var("P", ty.ptr<private_>(T));
+ });
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ b.Return(f2, b.Load(b.Access<ptr<private_, vec4<i32>, read_write>>(p, 3_u)));
+ });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<private_>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<private_>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<private_>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ b.Call(f0, P);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%b1 = block { # root
+ %P:ptr<private, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+}
+
+%f2 = func(%p:ptr<private, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b2 {
+ %b2 = block {
+ %4:ptr<private, vec4<i32>, read_write> = access %p, 3u
+ %5:vec4<i32> = load %4
+ ret %5
+ }
+}
+%f1 = func(%p_1:ptr<private, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_1: 'p'
+ %b3 = block {
+ %8:ptr<private, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %9:vec4<i32> = call %f2, %8
+ ret %9
+ }
+}
+%f0 = func(%p_2:ptr<private, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b4 { # %p_2: 'p'
+ %b4 = block {
+ %12:ptr<private, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %13:vec4<i32> = call %f1, %12
+ ret %13
+ }
+}
+%main = func():void -> %b5 {
+ %b5 = block {
+ %15:vec4<i32> = call %f0, %P
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace private_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -3841,6 +4415,567 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_DirectVariableAccessTest_FunctionAS, Enabled_CallChaining) {
+ auto* Inner =
+ ty.Struct(mod.symbols.New("Inner"), {
+ {mod.symbols.Register("mat"), ty.mat3x4<f32>()},
+ });
+ auto* Outer =
+ ty.Struct(mod.symbols.New("Outer"), {
+ {mod.symbols.Register("arr"), ty.array(Inner, 4)},
+ {mod.symbols.Register("mat"), ty.mat3x4<f32>()},
+ });
+
+ auto* f0 = b.Function("f0", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function, vec4<f32>>());
+ f0->SetParams({p});
+ b.Append(f0->Block(), [&] { b.Return(f0, b.LoadVectorElement(p, 0_u)); });
+ }
+
+ auto* f1 = b.Function("f1", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function, mat3x4<f32>>());
+ f1->SetParams({p});
+ b.Append(f1->Block(), [&] {
+ auto* res = b.Var<function, f32>("res");
+ {
+ // res += f0(&(*p)[1]);
+ auto* call_0 = b.Call(f0, b.Access(ty.ptr<function, vec4<f32>>(), p, 1_i));
+ b.Store(res, b.Add(ty.f32(), b.Load(res), call_0));
+ }
+ {
+ // let p_vec = &(*p)[1];
+ // res += f0(p_vec);
+ auto* p_vec = b.Access(ty.ptr<function, vec4<f32>>(), p, 1_i);
+ b.ir.SetName(p_vec, "p_vec");
+ auto* call_0 = b.Call(f0, p_vec);
+ b.Store(res, b.Add(ty.f32(), b.Load(res), call_0));
+ }
+ b.Return(f1, b.Load(res));
+ });
+ }
+
+ auto* f2 = b.Function("f2", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(Inner));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ auto* p_mat = b.Access(ty.ptr<function, mat3x4<f32>>(), p, 0_u);
+ b.ir.SetName(p_mat, "p_mat");
+ b.Return(f2, b.Call(f1, p_mat));
+ });
+ }
+
+ auto* f3 = b.Function("f3", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(ty.array(Inner, 4)));
+ f3->SetParams({p});
+ b.Append(f3->Block(), [&] {
+ auto* p_inner = b.Access(ty.ptr<function>(Inner), p, 3_i);
+ b.ir.SetName(p_inner, "p_inner");
+ b.Return(f3, b.Call(f2, p_inner));
+ });
+ }
+
+ auto* f4 = b.Function("f4", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(Outer));
+ f4->SetParams({p});
+ b.Append(f4->Block(), [&] {
+ auto* access = b.Access(ty.ptr<function>(ty.array(Inner, 4)), p, 0_u);
+ b.Return(f4, b.Call(f3, access));
+ });
+ }
+
+ auto* fn_b = b.Function("b", ty.void_());
+ b.Append(fn_b->Block(), [&] {
+ auto F = b.Var("F", ty.ptr<function>(Outer));
+ b.Call(f4, F);
+ b.Return(fn_b);
+ });
+
+ auto* src = R"(
+Inner = struct @align(16) {
+ mat:mat3x4<f32> @offset(0)
+}
+
+Outer = struct @align(16) {
+ arr:array<Inner, 4> @offset(0)
+ mat:mat3x4<f32> @offset(192)
+}
+
+%f0 = func(%p:ptr<function, vec4<f32>, read_write>):f32 -> %b1 {
+ %b1 = block {
+ %3:f32 = load_vector_element %p, 0u
+ ret %3
+ }
+}
+%f1 = func(%p_1:ptr<function, mat3x4<f32>, read_write>):f32 -> %b2 { # %p_1: 'p'
+ %b2 = block {
+ %res:ptr<function, f32, read_write> = var
+ %7:ptr<function, vec4<f32>, read_write> = access %p_1, 1i
+ %8:f32 = call %f0, %7
+ %9:f32 = load %res
+ %10:f32 = add %9, %8
+ store %res, %10
+ %p_vec:ptr<function, vec4<f32>, read_write> = access %p_1, 1i
+ %12:f32 = call %f0, %p_vec
+ %13:f32 = load %res
+ %14:f32 = add %13, %12
+ store %res, %14
+ %15:f32 = load %res
+ ret %15
+ }
+}
+%f2 = func(%p_2:ptr<function, Inner, read_write>):f32 -> %b3 { # %p_2: 'p'
+ %b3 = block {
+ %p_mat:ptr<function, mat3x4<f32>, read_write> = access %p_2, 0u
+ %19:f32 = call %f1, %p_mat
+ ret %19
+ }
+}
+%f3 = func(%p_3:ptr<function, array<Inner, 4>, read_write>):f32 -> %b4 { # %p_3: 'p'
+ %b4 = block {
+ %p_inner:ptr<function, Inner, read_write> = access %p_3, 3i
+ %23:f32 = call %f2, %p_inner
+ ret %23
+ }
+}
+%f4 = func(%p_4:ptr<function, Outer, read_write>):f32 -> %b5 { # %p_4: 'p'
+ %b5 = block {
+ %26:ptr<function, array<Inner, 4>, read_write> = access %p_4, 0u
+ %27:f32 = call %f3, %26
+ ret %27
+ }
+}
+%b = func():void -> %b6 {
+ %b6 = block {
+ %F:ptr<function, Outer, read_write> = var
+ %30:f32 = call %f4, %F
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Inner = struct @align(16) {
+ mat:mat3x4<f32> @offset(0)
+}
+
+Outer = struct @align(16) {
+ arr:array<Inner, 4> @offset(0)
+ mat:mat3x4<f32> @offset(192)
+}
+
+%f0_P_arr_X_mat_X = func(%p_root:ptr<function, Outer, read_write>, %p_indices:array<u32, 2>):f32 -> %b1 {
+ %b1 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<function, vec4<f32>, read_write> = access %p_root, 0u, %4, 0u, %5
+ %7:f32 = load_vector_element %6, 0u
+ ret %7
+ }
+}
+%f1_P_arr_X_mat = func(%p_root_1:ptr<function, Outer, read_write>, %p_indices_1:array<u32, 1>):f32 -> %b2 { # %p_root_1: 'p_root', %p_indices_1: 'p_indices'
+ %b2 = block {
+ %11:u32 = access %p_indices_1, 0u
+ %res:ptr<function, f32, read_write> = var
+ %13:u32 = convert 1i
+ %14:array<u32, 2> = construct %11, %13
+ %15:f32 = call %f0_P_arr_X_mat_X, %p_root_1, %14
+ %16:f32 = load %res
+ %17:f32 = add %16, %15
+ store %res, %17
+ %18:u32 = convert 1i
+ %19:array<u32, 2> = construct %11, %18
+ %20:f32 = call %f0_P_arr_X_mat_X, %p_root_1, %19
+ %21:f32 = load %res
+ %22:f32 = add %21, %20
+ store %res, %22
+ %23:f32 = load %res
+ ret %23
+ }
+}
+%f2_P_arr_X = func(%p_root_2:ptr<function, Outer, read_write>, %p_indices_2:array<u32, 1>):f32 -> %b3 { # %p_root_2: 'p_root', %p_indices_2: 'p_indices'
+ %b3 = block {
+ %27:u32 = access %p_indices_2, 0u
+ %28:array<u32, 1> = construct %27
+ %29:f32 = call %f1_P_arr_X_mat, %p_root_2, %28
+ ret %29
+ }
+}
+%f3_P_arr = func(%p_root_3:ptr<function, Outer, read_write>):f32 -> %b4 { # %p_root_3: 'p_root'
+ %b4 = block {
+ %32:u32 = convert 3i
+ %33:array<u32, 1> = construct %32
+ %34:f32 = call %f2_P_arr_X, %p_root_3, %33
+ ret %34
+ }
+}
+%f4_P = func(%p_root_4:ptr<function, Outer, read_write>):f32 -> %b5 { # %p_root_4: 'p_root'
+ %b5 = block {
+ %37:f32 = call %f3_P_arr, %p_root_4
+ ret %37
+ }
+}
+%b = func():void -> %b6 {
+ %b6 = block {
+ %F:ptr<function, Outer, read_write> = var
+ %40:f32 = call %f4_P, %F
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, kTransformFunction);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_DirectVariableAccessTest_FunctionAS, Disabled_CallChaining) {
+ auto* Inner =
+ ty.Struct(mod.symbols.New("Inner"), {
+ {mod.symbols.Register("mat"), ty.mat3x4<f32>()},
+ });
+ auto* Outer =
+ ty.Struct(mod.symbols.New("Outer"), {
+ {mod.symbols.Register("arr"), ty.array(Inner, 4)},
+ {mod.symbols.Register("mat"), ty.mat3x4<f32>()},
+ });
+
+ auto* f0 = b.Function("f0", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function, vec4<f32>>());
+ f0->SetParams({p});
+ b.Append(f0->Block(), [&] { b.Return(f0, b.LoadVectorElement(p, 0_u)); });
+ }
+
+ auto* f1 = b.Function("f1", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function, mat3x4<f32>>());
+ f1->SetParams({p});
+ b.Append(f1->Block(), [&] {
+ auto* res = b.Var<function, f32>("res");
+ {
+ // res += f0(&(*p)[1]);
+ auto* call_0 = b.Call(f0, b.Access(ty.ptr<function, vec4<f32>>(), p, 1_i));
+ b.Store(res, b.Add(ty.f32(), b.Load(res), call_0));
+ }
+ {
+ // let p_vec = &(*p)[1];
+ // res += f0(p_vec);
+ auto* p_vec = b.Access(ty.ptr<function, vec4<f32>>(), p, 1_i);
+ b.ir.SetName(p_vec, "p_vec");
+ auto* call_0 = b.Call(f0, p_vec);
+ b.Store(res, b.Add(ty.f32(), b.Load(res), call_0));
+ }
+ b.Return(f1, b.Load(res));
+ });
+ }
+
+ auto* f2 = b.Function("f2", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(Inner));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ auto* p_mat = b.Access(ty.ptr<function, mat3x4<f32>>(), p, 0_u);
+ b.ir.SetName(p_mat, "p_mat");
+ b.Return(f2, b.Call(f1, p_mat));
+ });
+ }
+
+ auto* f3 = b.Function("f3", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(ty.array(Inner, 4)));
+ f3->SetParams({p});
+ b.Append(f3->Block(), [&] {
+ auto* p_inner = b.Access(ty.ptr<function>(Inner), p, 3_i);
+ b.ir.SetName(p_inner, "p_inner");
+ b.Return(f3, b.Call(f2, p_inner));
+ });
+ }
+
+ auto* f4 = b.Function("f4", ty.f32());
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(Outer));
+ f4->SetParams({p});
+ b.Append(f4->Block(), [&] {
+ auto* access = b.Access(ty.ptr<function>(ty.array(Inner, 4)), p, 0_u);
+ b.Return(f4, b.Call(f3, access));
+ });
+ }
+
+ auto* fn_b = b.Function("b", ty.void_());
+ b.Append(fn_b->Block(), [&] {
+ auto F = b.Var("F", ty.ptr<function>(Outer));
+ b.Call(f4, F);
+ b.Return(fn_b);
+ });
+
+ auto* src = R"(
+Inner = struct @align(16) {
+ mat:mat3x4<f32> @offset(0)
+}
+
+Outer = struct @align(16) {
+ arr:array<Inner, 4> @offset(0)
+ mat:mat3x4<f32> @offset(192)
+}
+
+%f0 = func(%p:ptr<function, vec4<f32>, read_write>):f32 -> %b1 {
+ %b1 = block {
+ %3:f32 = load_vector_element %p, 0u
+ ret %3
+ }
+}
+%f1 = func(%p_1:ptr<function, mat3x4<f32>, read_write>):f32 -> %b2 { # %p_1: 'p'
+ %b2 = block {
+ %res:ptr<function, f32, read_write> = var
+ %7:ptr<function, vec4<f32>, read_write> = access %p_1, 1i
+ %8:f32 = call %f0, %7
+ %9:f32 = load %res
+ %10:f32 = add %9, %8
+ store %res, %10
+ %p_vec:ptr<function, vec4<f32>, read_write> = access %p_1, 1i
+ %12:f32 = call %f0, %p_vec
+ %13:f32 = load %res
+ %14:f32 = add %13, %12
+ store %res, %14
+ %15:f32 = load %res
+ ret %15
+ }
+}
+%f2 = func(%p_2:ptr<function, Inner, read_write>):f32 -> %b3 { # %p_2: 'p'
+ %b3 = block {
+ %p_mat:ptr<function, mat3x4<f32>, read_write> = access %p_2, 0u
+ %19:f32 = call %f1, %p_mat
+ ret %19
+ }
+}
+%f3 = func(%p_3:ptr<function, array<Inner, 4>, read_write>):f32 -> %b4 { # %p_3: 'p'
+ %b4 = block {
+ %p_inner:ptr<function, Inner, read_write> = access %p_3, 3i
+ %23:f32 = call %f2, %p_inner
+ ret %23
+ }
+}
+%f4 = func(%p_4:ptr<function, Outer, read_write>):f32 -> %b5 { # %p_4: 'p'
+ %b5 = block {
+ %26:ptr<function, array<Inner, 4>, read_write> = access %p_4, 0u
+ %27:f32 = call %f3, %26
+ ret %27
+ }
+}
+%b = func():void -> %b6 {
+ %b6 = block {
+ %F:ptr<function, Outer, read_write> = var
+ %30:f32 = call %f4, %F
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_DirectVariableAccessTest_FunctionAS, Enabled_CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ b.Return(f2, b.Load(b.Access<ptr<function, vec4<i32>, read_write>>(p, 3_u)));
+ });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<function>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<function>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ auto* F = b.Var("F", ty.ptr<function>(T));
+ b.Call(f0, F);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%f2 = func(%p:ptr<function, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b1 {
+ %b1 = block {
+ %3:ptr<function, vec4<i32>, read_write> = access %p, 3u
+ %4:vec4<i32> = load %3
+ ret %4
+ }
+}
+%f1 = func(%p_1:ptr<function, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b2 { # %p_1: 'p'
+ %b2 = block {
+ %7:ptr<function, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %8:vec4<i32> = call %f2, %7
+ ret %8
+ }
+}
+%f0 = func(%p_2:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_2: 'p'
+ %b3 = block {
+ %11:ptr<function, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %12:vec4<i32> = call %f1, %11
+ ret %12
+ }
+}
+%main = func():void -> %b4 {
+ %b4 = block {
+ %F:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+ %15:vec4<i32> = call %f0, %F
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%f2_P_X_X = func(%p_root:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>, %p_indices:array<u32, 2>):vec4<i32> -> %b1 {
+ %b1 = block {
+ %4:u32 = access %p_indices, 0u
+ %5:u32 = access %p_indices, 1u
+ %6:ptr<function, array<vec4<i32>, 5>, read_write> = access %p_root, %4, %5
+ %7:ptr<function, vec4<i32>, read_write> = access %6, 3u
+ %8:vec4<i32> = load %7
+ ret %8
+ }
+}
+%f1_P_X = func(%p_root_1:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>, %p_indices_1:array<u32, 1>):vec4<i32> -> %b2 { # %p_root_1: 'p_root', %p_indices_1: 'p_indices'
+ %b2 = block {
+ %12:u32 = access %p_indices_1, 0u
+ %13:array<u32, 2> = construct %12, 2u
+ %14:vec4<i32> = call %f2_P_X_X, %p_root_1, %13
+ ret %14
+ }
+}
+%f0_P = func(%p_root_2:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_root_2: 'p_root'
+ %b3 = block {
+ %17:array<u32, 1> = construct 1u
+ %18:vec4<i32> = call %f1_P_X, %p_root_2, %17
+ ret %18
+ }
+}
+%main = func():void -> %b4 {
+ %b4 = block {
+ %F:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+ %21:vec4<i32> = call %f0_P, %F
+ ret
+ }
+}
+)";
+
+ Run(DirectVariableAccess, kTransformFunction);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_DirectVariableAccessTest_FunctionAS, Disabled_CallChaining2) {
+ auto* T3 = ty.vec4<i32>();
+ auto* T2 = ty.array(T3, 5);
+ auto* T1 = ty.array(T2, 5);
+ auto* T = ty.array(T1, 5);
+
+ auto* f2 = b.Function("f2", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T2));
+ f2->SetParams({p});
+ b.Append(f2->Block(), [&] {
+ b.Return(f2, b.Load(b.Access<ptr<function, vec4<i32>, read_write>>(p, 3_u)));
+ });
+ }
+
+ auto* f1 = b.Function("f1", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T1));
+ f1->SetParams({p});
+ b.Append(f1->Block(),
+ [&] { b.Return(f1, b.Call(f2, b.Access(ty.ptr<function>(T2), p, 2_u))); });
+ }
+
+ auto* f0 = b.Function("f0", T3);
+ {
+ auto* p = b.FunctionParam("p", ty.ptr<function>(T));
+ f0->SetParams({p});
+ b.Append(f0->Block(),
+ [&] { b.Return(f0, b.Call(f1, b.Access(ty.ptr<function>(T1), p, 1_u))); });
+ }
+
+ auto* main = b.Function("main", ty.void_());
+ b.Append(main->Block(), [&] {
+ auto* F = b.Var("F", ty.ptr<function>(T));
+ b.Call(f0, F);
+ b.Return(main);
+ });
+
+ auto* src = R"(
+%f2 = func(%p:ptr<function, array<vec4<i32>, 5>, read_write>):vec4<i32> -> %b1 {
+ %b1 = block {
+ %3:ptr<function, vec4<i32>, read_write> = access %p, 3u
+ %4:vec4<i32> = load %3
+ ret %4
+ }
+}
+%f1 = func(%p_1:ptr<function, array<array<vec4<i32>, 5>, 5>, read_write>):vec4<i32> -> %b2 { # %p_1: 'p'
+ %b2 = block {
+ %7:ptr<function, array<vec4<i32>, 5>, read_write> = access %p_1, 2u
+ %8:vec4<i32> = call %f2, %7
+ ret %8
+ }
+}
+%f0 = func(%p_2:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write>):vec4<i32> -> %b3 { # %p_2: 'p'
+ %b3 = block {
+ %11:ptr<function, array<array<vec4<i32>, 5>, 5>, read_write> = access %p_2, 1u
+ %12:vec4<i32> = call %f1, %11
+ ret %12
+ }
+}
+%main = func():void -> %b4 {
+ %b4 = block {
+ %F:ptr<function, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
+ %15:vec4<i32> = call %f0, %F
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(DirectVariableAccess, DirectVariableAccessOptions{});
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace function_as_tests
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 5546f2c..2aa16f8 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -598,6 +598,12 @@
AddError(call, UserCall::kFunctionOperandOffset,
InstError(call, "call target is not part of the module"));
}
+
+ if (call->Target()->Stage() != Function::PipelineStage::kUndefined) {
+ AddError(call, UserCall::kFunctionOperandOffset,
+ InstError(call, "call target must not have a pipeline stage"));
+ }
+
auto args = call->Args();
auto params = call->Target()->Params();
if (args.Length() != params.Length()) {
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index d81f679..3c2733a 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -179,6 +179,42 @@
)");
}
+TEST_F(IR_ValidatorTest, CallToEntryPointFunction) {
+ auto* f = b.Function("f", ty.void_());
+ auto* g = b.Function("g", ty.void_(), Function::PipelineStage::kCompute);
+
+ b.Append(f->Block(), [&] {
+ b.Call(g);
+ b.Return(f);
+ });
+ b.Append(g->Block(), [&] { b.Return(g); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.str(),
+ R"(:3:20 error: call: call target must not have a pipeline stage
+ %2:void = call %g
+ ^^
+
+:2:3 note: In block
+ %b1 = block {
+ ^^^^^^^^^^^
+
+note: # Disassembly
+%f = func():void -> %b1 {
+ %b1 = block {
+ %2:void = call %g
+ ret
+ }
+}
+%g = @compute func():void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, CallToFunctionTooFewArguments) {
auto* g = b.Function("g", ty.void_());
g->SetParams({b.FunctionParam<i32>(), b.FunctionParam<i32>()});
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc b/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
index df9df7d..c73dead 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
@@ -84,16 +84,13 @@
/// DynamicIndex is used by DirectVariableAccess::State::AccessOp to indicate an array, matrix or
/// vector index.
struct DynamicIndex {
- /// The index of the expression in DirectVariableAccess::State::AccessChain::dynamic_indices
- size_t slot = 0;
-
/// @return a hash code for this object
- size_t HashCode() const { return Hash(slot); }
+ size_t HashCode() const { return 42 /* empty struct: any number will do */; }
};
/// Inequality operator for DynamicIndex
-bool operator!=(const DynamicIndex& a, const DynamicIndex& b) {
- return a.slot != b.slot;
+bool operator!=(const DynamicIndex&, const DynamicIndex&) {
+ return false; // empty struct: two DynamicIndex objects are always equal
}
/// AccessOp describes a single access in an access chain.
@@ -484,7 +481,7 @@
// Store the index expression into AccessChain::dynamic_indices, append a
// DynamicIndex to the chain, and move the chain to the index accessor expression.
if (auto* chain = take_chain(a->Object())) {
- chain->ops.Push(DynamicIndex{chain->dynamic_indices.Length()});
+ chain->ops.Push(DynamicIndex{});
chain->dynamic_indices.Push(a->Index());
}
},
@@ -1025,20 +1022,23 @@
// Chain starts with a pointer parameter.
// Replace this with the variant's incoming shape. This will bring the expression up to
// the incoming pointer.
+ size_t next_dyn_idx_from_indices = 0;
auto indices =
clone_state->current_variant->ptr_param_symbols.Find(root_param)->indices;
for (auto param_access : incoming_shape->ops) {
- chain_expr = BuildAccessExpr(chain_expr, param_access, [&](size_t i) {
- return b.IndexAccessor(indices, AInt(i));
+ chain_expr = BuildAccessExpr(chain_expr, param_access, [&] {
+ return b.IndexAccessor(indices, AInt(next_dyn_idx_from_indices++));
});
}
// Now build the expression chain within the function.
// For each access in the chain (excluding the pointer parameter)...
+ size_t next_dyn_idx_from_chain = 0;
for (auto& op : chain->ops) {
- chain_expr = BuildAccessExpr(chain_expr, op, [&](size_t i) {
- return BuildDynamicIndex(chain->dynamic_indices[i], false);
+ chain_expr = BuildAccessExpr(chain_expr, op, [&] {
+ return BuildDynamicIndex(chain->dynamic_indices[next_dyn_idx_from_chain++],
+ false);
});
}
@@ -1131,13 +1131,13 @@
/// The returned expression will always be of a reference type.
/// @param expr the input expression
/// @param access the access to perform on the current expression
- /// @param dynamic_index a function that obtains the i'th dynamic index
+ /// @param dynamic_index a function that obtains the next dynamic index
const Expression* BuildAccessExpr(const Expression* expr,
const AccessOp& access,
- std::function<const Expression*(size_t)> dynamic_index) {
- if (auto* dyn_idx = std::get_if<DynamicIndex>(&access)) {
+ std::function<const Expression*()> dynamic_index) {
+ if (std::holds_alternative<DynamicIndex>(access)) {
/// The access uses a dynamic (runtime-expression) index.
- auto* idx = dynamic_index(dyn_idx->slot);
+ auto* idx = dynamic_index();
return b.IndexAccessor(expr, idx);
}
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
index cb5598f..10f6421 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
@@ -1035,12 +1035,6 @@
alias U_arr_X_mat_X = array<u32, 2u>;
fn f0_U_arr_X_mat_X(p : U_arr_X_mat_X) -> f32 {
- return U.arr[p[0]].mat[p[0]].x;
-}
-
-alias U_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_U_arr_X_mat_X_1(p : U_arr_X_mat_X_1) -> f32 {
return U.arr[p[0]].mat[p[1]].x;
}
@@ -1054,11 +1048,11 @@
res += f0_U_mat_X(U_mat_X(1));
}
{
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
{
let p_vec = &(U.arr[2].mat[1]);
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1075,11 +1069,11 @@
res += f0_U_arr_X_mat_X(U_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
{
let p_vec = &(U.arr[2].mat[1]);
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1198,12 +1192,6 @@
alias U_arr_X_mat_X = array<u32, 2u>;
fn f0_U_arr_X_mat_X(p : U_arr_X_mat_X) -> f32 {
- return (&(U.arr[p[0]].mat[p[0]])).x;
-}
-
-alias U_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_U_arr_X_mat_X_1(p : U_arr_X_mat_X_1) -> f32 {
return (&(U.arr[p[0]].mat[p[1]])).x;
}
@@ -1217,11 +1205,11 @@
res += f0_U_mat_X(U_mat_X(1));
}
{
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
{
let p_vec = &(U.arr[2].mat[1]);
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1238,11 +1226,11 @@
res += f0_U_arr_X_mat_X(U_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
{
let p_vec = &(U.arr[2].mat[1]);
- res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1273,6 +1261,76 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessUniformASTest, CallChaining2) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+alias T2 = array<T3, 5>;
+alias T1 = array<T2, 5>;
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<uniform> input : T;
+
+fn f2(p : ptr<uniform, T2>) -> T3 {
+ return (*p)[3];
+}
+
+fn f1(p : ptr<uniform, T1>) -> T3 {
+ return f2(&(*p)[2]);
+}
+
+fn f0(p : ptr<uniform, T>) -> T3 {
+ return f1(&(*p)[1]);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ f0(&input);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+
+alias T2 = array<T3, 5>;
+
+alias T1 = array<T2, 5>;
+
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<uniform> input : T;
+
+alias input_X_X = array<u32, 2u>;
+
+fn f2_input_X_X(p : input_X_X) -> T3 {
+ return input[p[0]][p[1]][3];
+}
+
+alias input_X = array<u32, 1u>;
+
+fn f1_input_X(p : input_X) -> T3 {
+ return f2_input_X_X(input_X_X(p[0u], 2));
+}
+
+fn f0_input() -> T3 {
+ return f1_input_X(input_X(1));
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ f0_input();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace uniform_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -1492,12 +1550,6 @@
alias S_arr_X_mat_X = array<u32, 2u>;
fn f0_S_arr_X_mat_X(p : S_arr_X_mat_X) -> f32 {
- return S.arr[p[0]].mat[p[0]].x;
-}
-
-alias S_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_S_arr_X_mat_X_1(p : S_arr_X_mat_X_1) -> f32 {
return S.arr[p[0]].mat[p[1]].x;
}
@@ -1511,11 +1563,11 @@
res += f0_S_mat_X(S_mat_X(1));
}
{
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
{
let p_vec = &(S.arr[2].mat[1]);
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1532,11 +1584,11 @@
res += f0_S_arr_X_mat_X(S_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
{
let p_vec = &(S.arr[2].mat[1]);
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1655,12 +1707,6 @@
alias S_arr_X_mat_X = array<u32, 2u>;
fn f0_S_arr_X_mat_X(p : S_arr_X_mat_X) -> f32 {
- return (&(S.arr[p[0]].mat[p[0]])).x;
-}
-
-alias S_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_S_arr_X_mat_X_1(p : S_arr_X_mat_X_1) -> f32 {
return (&(S.arr[p[0]].mat[p[1]])).x;
}
@@ -1674,11 +1720,11 @@
res += f0_S_mat_X(S_mat_X(1));
}
{
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
{
let p_vec = &(S.arr[2].mat[1]);
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1695,11 +1741,11 @@
res += f0_S_arr_X_mat_X(S_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
{
let p_vec = &(S.arr[2].mat[1]);
- res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1730,6 +1776,76 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessStorageASTest, CallChaining2) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+alias T2 = array<T3, 5>;
+alias T1 = array<T2, 5>;
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<storage> input : T;
+
+fn f2(p : ptr<storage, T2>) -> T3 {
+ return (*p)[3];
+}
+
+fn f1(p : ptr<storage, T1>) -> T3 {
+ return f2(&(*p)[2]);
+}
+
+fn f0(p : ptr<storage, T>) -> T3 {
+ return f1(&(*p)[1]);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ f0(&input);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+
+alias T2 = array<T3, 5>;
+
+alias T1 = array<T2, 5>;
+
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<storage> input : T;
+
+alias input_X_X = array<u32, 2u>;
+
+fn f2_input_X_X(p : input_X_X) -> T3 {
+ return input[p[0]][p[1]][3];
+}
+
+alias input_X = array<u32, 1u>;
+
+fn f1_input_X(p : input_X) -> T3 {
+ return f2_input_X_X(input_X_X(p[0u], 2));
+}
+
+fn f0_input() -> T3 {
+ return f1_input_X(input_X(1));
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ f0_input();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace storage_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -1899,12 +2015,6 @@
alias W_arr_X_mat_X = array<u32, 2u>;
fn f0_W_arr_X_mat_X(p : W_arr_X_mat_X) -> f32 {
- return W.arr[p[0]].mat[p[0]].x;
-}
-
-alias W_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_W_arr_X_mat_X_1(p : W_arr_X_mat_X_1) -> f32 {
return W.arr[p[0]].mat[p[1]].x;
}
@@ -1918,11 +2028,11 @@
res += f0_W_mat_X(W_mat_X(1));
}
{
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
{
let p_vec = &(W.arr[2].mat[1]);
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
return res;
}
@@ -1939,11 +2049,11 @@
res += f0_W_arr_X_mat_X(W_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
{
let p_vec = &(W.arr[2].mat[1]);
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
return res;
}
@@ -2062,12 +2172,6 @@
alias W_arr_X_mat_X = array<u32, 2u>;
fn f0_W_arr_X_mat_X(p : W_arr_X_mat_X) -> f32 {
- return (&(W.arr[p[0]].mat[p[0]])).x;
-}
-
-alias W_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_W_arr_X_mat_X_1(p : W_arr_X_mat_X_1) -> f32 {
return (&(W.arr[p[0]].mat[p[1]])).x;
}
@@ -2081,11 +2185,11 @@
res += f0_W_mat_X(W_mat_X(1));
}
{
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
{
let p_vec = &(W.arr[2].mat[1]);
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
return res;
}
@@ -2102,11 +2206,11 @@
res += f0_W_arr_X_mat_X(W_arr_X_mat_X(p[0u], 1));
}
{
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
{
let p_vec = &(W.arr[2].mat[1]);
- res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(2, 1));
}
return res;
}
@@ -2137,6 +2241,81 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessWorkgroupASTest, CallChaining2) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+alias T2 = array<T3, 5>;
+alias T1 = array<T2, 5>;
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<storage, read> input : T;
+
+fn f2(p : ptr<workgroup, T2>) -> T3 {
+ return (*p)[3];
+}
+
+fn f1(p : ptr<workgroup, T1>) -> T3 {
+ return f2(&(*p)[2]);
+}
+
+fn f0(p : ptr<workgroup, T>) -> T3 {
+ return f1(&(*p)[1]);
+}
+
+var<workgroup> W : T;
+@compute @workgroup_size(1)
+fn main() {
+ W = input;
+ f0(&W);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+
+alias T2 = array<T3, 5>;
+
+alias T1 = array<T2, 5>;
+
+alias T = array<T1, 5>;
+
+@binding(0) @group(0) var<storage, read> input : T;
+
+alias W_X_X = array<u32, 2u>;
+
+fn f2_W_X_X(p : W_X_X) -> T3 {
+ return W[p[0]][p[1]][3];
+}
+
+alias W_X = array<u32, 1u>;
+
+fn f1_W_X(p : W_X) -> T3 {
+ return f2_W_X_X(W_X_X(p[0u], 2));
+}
+
+fn f0_W() -> T3 {
+ return f1_W_X(W_X(1));
+}
+
+var<workgroup> W : T;
+
+@compute @workgroup_size(1)
+fn main() {
+ W = input;
+ f0_W();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace workgroup_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -2534,18 +2713,12 @@
alias F_arr_X_mat_X = array<u32, 2u>;
fn f0_F_arr_X_mat_X(p_base : ptr<private, Outer>, p_indices : F_arr_X_mat_X) -> f32 {
- return (*(p_base)).arr[p_indices[0]].mat[p_indices[0]].x;
-}
-
-alias F_arr_X_mat_X_1 = array<u32, 2u>;
-
-fn f0_F_arr_X_mat_X_1(p_base : ptr<private, Outer>, p_indices : F_arr_X_mat_X_1) -> f32 {
return (*(p_base)).arr[p_indices[0]].mat[p_indices[1]].x;
}
alias F_mat_X_1 = array<u32, 1u>;
-alias F_arr_X_mat_X_2 = array<u32, 2u>;
+alias F_arr_X_mat_X_1 = array<u32, 2u>;
fn f1_F_mat(p : ptr<private, Outer>) -> f32 {
var res : f32;
@@ -2557,34 +2730,34 @@
res += f0_F_mat_X(p, F_mat_X_1(1));
}
{
- res += f0_F_arr_X_mat_X_1(&(P), F_arr_X_mat_X_2(2, 1));
+ res += f0_F_arr_X_mat_X(&(P), F_arr_X_mat_X_1(2, 1));
}
{
let p_vec = &(P.arr[2].mat[1]);
- res += f0_F_arr_X_mat_X_1(&(P), F_arr_X_mat_X_2(2, 1));
+ res += f0_F_arr_X_mat_X(&(P), F_arr_X_mat_X_1(2, 1));
}
return res;
}
alias F_arr_X_mat = array<u32, 1u>;
-alias F_arr_X_mat_X_3 = array<u32, 2u>;
+alias F_arr_X_mat_X_2 = array<u32, 2u>;
fn f1_F_arr_X_mat(p_base : ptr<private, Outer>, p_indices : F_arr_X_mat) -> f32 {
var res : f32;
{
- res += f0_F_arr_X_mat_X(p_base, F_arr_X_mat_X_3(p_indices[0u], 1));
+ res += f0_F_arr_X_mat_X(p_base, F_arr_X_mat_X_2(p_indices[0u], 1));
}
{
let p_vec = &((*(p_base)).arr[p_indices[0]].mat[1]);
- res += f0_F_arr_X_mat_X(p_base, F_arr_X_mat_X_3(p_indices[0u], 1));
+ res += f0_F_arr_X_mat_X(p_base, F_arr_X_mat_X_2(p_indices[0u], 1));
}
{
- res += f0_F_arr_X_mat_X_1(&(P), F_arr_X_mat_X_2(2, 1));
+ res += f0_F_arr_X_mat_X(&(P), F_arr_X_mat_X_1(2, 1));
}
{
let p_vec = &(P.arr[2].mat[1]);
- res += f0_F_arr_X_mat_X_1(&(P), F_arr_X_mat_X_2(2, 1));
+ res += f0_F_arr_X_mat_X(&(P), F_arr_X_mat_X_1(2, 1));
}
return res;
}
@@ -2619,6 +2792,80 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessPrivateASTest, Enabled_CallChaining2) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+alias T2 = array<T3, 5>;
+alias T1 = array<T2, 5>;
+alias T = array<T1, 5>;
+
+fn f2(p : ptr<private, T2>) -> T3 {
+ return (*p)[3];
+}
+
+fn f1(p : ptr<private, T1>) -> T3 {
+ return f2(&(*p)[2]);
+}
+
+fn f0(p : ptr<private, T>) -> T3 {
+ return f1(&(*p)[1]);
+}
+
+var<private> P : T;
+
+@compute @workgroup_size(1)
+fn main() {
+ f0(&P);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+
+alias T2 = array<T3, 5>;
+
+alias T1 = array<T2, 5>;
+
+alias T = array<T1, 5>;
+
+alias F_X_X = array<u32, 2u>;
+
+fn f2_F_X_X(p_base : ptr<private, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>, p_indices : F_X_X) -> T3 {
+ return (*(p_base))[p_indices[0]][p_indices[1]][3];
+}
+
+alias F_X = array<u32, 1u>;
+
+alias F_X_X_1 = array<u32, 2u>;
+
+fn f1_F_X(p_base : ptr<private, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>, p_indices : F_X) -> T3 {
+ return f2_F_X_X(p_base, F_X_X_1(p_indices[0u], 2));
+}
+
+alias F_X_1 = array<u32, 1u>;
+
+fn f0_F(p : ptr<private, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>) -> T3 {
+ return f1_F_X(p, F_X_1(1));
+}
+
+var<private> P : T;
+
+@compute @workgroup_size(1)
+fn main() {
+ f0_F(&(P));
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src, EnablePrivate());
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DirectVariableAccessPrivateASTest, Disabled_CallChaining) {
auto* src = R"(
enable chromium_experimental_full_ptr_parameters;
@@ -2917,6 +3164,187 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessFunctionASTest, Enabled_CallChaining) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+};
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+};
+
+fn f0(p : ptr<function, vec4<f32>>) -> f32 {
+ return (*p).x;
+}
+
+fn f1(p : ptr<function, mat3x4<f32>>) -> f32 {
+ var res : f32;
+ {
+ // call f0() with inline usage of p
+ res += f0(&(*p)[1]);
+ }
+ {
+ // call f0() with pointer-let usage of p
+ let p_vec = &(*p)[1];
+ res += f0(p_vec);
+ }
+ return res;
+}
+
+fn f2(p : ptr<function, Inner>) -> f32 {
+ let p_mat = &(*p).mat;
+ return f1(p_mat);
+}
+
+fn f3(p : ptr<function, InnerArr>) -> f32 {
+ let p_inner = &(*p)[3];
+ return f2(p_inner);
+}
+
+fn f4(p : ptr<function, Outer>) -> f32 {
+ return f3(&(*p).arr);
+}
+
+fn b() {
+ var S : Outer;
+ f4(&S);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+}
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+}
+
+fn f0(p : ptr<function, vec4<f32>>) -> f32 {
+ return (*(p)).x;
+}
+
+fn f1(p : ptr<function, mat3x4<f32>>) -> f32 {
+ var res : f32;
+ {
+ res += f0(&((*(p))[1]));
+ }
+ {
+ let p_vec = &((*(p))[1]);
+ res += f0(p_vec);
+ }
+ return res;
+}
+
+fn f2(p : ptr<function, Inner>) -> f32 {
+ let p_mat = &((*(p)).mat);
+ return f1(p_mat);
+}
+
+fn f3(p : ptr<function, InnerArr>) -> f32 {
+ let p_inner = &((*(p))[3]);
+ return f2(p_inner);
+}
+
+fn f4(p : ptr<function, Outer>) -> f32 {
+ return f3(&((*(p)).arr));
+}
+
+fn b() {
+ var S : Outer;
+ f4(&(S));
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DirectVariableAccessFunctionASTest, Enabled_CallChaining2) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+alias T2 = array<T3, 5>;
+alias T1 = array<T2, 5>;
+alias T = array<T1, 5>;
+
+fn f2(p : ptr<function, T2>) -> T3 {
+ return (*p)[3];
+}
+
+fn f1(p : ptr<function, T1>) -> T3 {
+ return f2(&(*p)[2]);
+}
+
+fn f0(p : ptr<function, T>) -> T3 {
+ return f1(&(*p)[1]);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var v : T;
+ f0(&v);
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+alias T3 = vec4i;
+
+alias T2 = array<T3, 5>;
+
+alias T1 = array<T2, 5>;
+
+alias T = array<T1, 5>;
+
+alias F_X_X = array<u32, 2u>;
+
+fn f2_F_X_X(p_base : ptr<function, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>, p_indices : F_X_X) -> T3 {
+ return (*(p_base))[p_indices[0]][p_indices[1]][3];
+}
+
+alias F_X = array<u32, 1u>;
+
+alias F_X_X_1 = array<u32, 2u>;
+
+fn f1_F_X(p_base : ptr<function, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>, p_indices : F_X) -> T3 {
+ return f2_F_X_X(p_base, F_X_X_1(p_indices[0u], 2));
+}
+
+alias F_X_1 = array<u32, 1u>;
+
+fn f0_F(p : ptr<function, array<array<array<vec4<i32>, 5u>, 5u>, 5u>>) -> T3 {
+ return f1_F_X(p, F_X_1(1));
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var v : T;
+ f0_F(&(v));
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src, EnableFunction());
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DirectVariableAccessFunctionASTest, Disabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
enable chromium_experimental_full_ptr_parameters;
@@ -3249,7 +3677,7 @@
alias S_X_X_X_X = array<u32, 4u>;
fn a_S_X_X_X_X(pre : i32, i : S_X_X_X_X, post : i32) -> i32 {
- return S[i[0]][i[0]][i[1]][i[2]];
+ return S[i[0]][i[1]][i[2]][i[3]];
}
alias S_X = array<u32, 1u>;
diff --git a/src/tint/lang/wgsl/reader/parser/additive_expression_test.cc b/src/tint/lang/wgsl/reader/parser/additive_expression_test.cc
index 782d922..0317b9a 100644
--- a/src/tint/lang/wgsl/reader/parser/additive_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/additive_expression_test.cc
@@ -33,15 +33,15 @@
TEST_F(WGSLParserTest, AdditiveExpression_Parses_Plus) {
auto p = parser("a + b");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 6u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -59,7 +59,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_Parses_Minus) {
auto p = parser("a - b");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -80,7 +80,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_Parses_MinusMinus) {
auto p = parser("a--b");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -105,7 +105,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_Parses_MultipleOps) {
auto p = parser("a - b + c - d");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -151,7 +151,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_Parses_MultipleOps_MixedMultiplication) {
auto p = parser("a - b * c - d");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -197,7 +197,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_InvalidRHS) {
auto p = parser("a + if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
EXPECT_TRUE(p->has_error());
@@ -207,7 +207,7 @@
TEST_F(WGSLParserTest, AdditiveExpression_NoMatch_ReturnsLHS) {
auto p = parser("a true");
auto lhs = p->unary_expression();
- auto e = p->expect_additive_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_additive_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/bitwise_expression_test.cc b/src/tint/lang/wgsl/reader/parser/bitwise_expression_test.cc
index 2a3fc71..4660dff 100644
--- a/src/tint/lang/wgsl/reader/parser/bitwise_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/bitwise_expression_test.cc
@@ -33,7 +33,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_NoOp) {
auto p = parser("a true");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
@@ -43,16 +43,16 @@
TEST_F(WGSLParserTest, BitwiseExpr_Or_Parses) {
auto p = parser("a | true");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 9u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -69,7 +69,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_Or_Parses_Multiple) {
auto p = parser("a | true | b");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
@@ -103,7 +103,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_Or_InvalidRHS) {
auto p = parser("true | if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.matched);
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
@@ -114,16 +114,16 @@
TEST_F(WGSLParserTest, BitwiseExpr_Xor_Parses) {
auto p = parser("a ^ true");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 9u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -140,7 +140,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_Xor_Parses_Multiple) {
auto p = parser("a ^ true ^ b");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
@@ -173,7 +173,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_Xor_InvalidRHS) {
auto p = parser("true ^ if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.matched);
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
@@ -184,16 +184,16 @@
TEST_F(WGSLParserTest, BitwiseExpr_And_Parses) {
auto p = parser("a & true");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 9u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -210,7 +210,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_And_Parses_Multiple) {
auto p = parser("a & true & b");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.matched);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
@@ -243,7 +243,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_And_Parses_AndAnd) {
auto p = parser("a & true &&b");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
// bitwise_expression_post_unary_expression returns before parsing '&&'
EXPECT_TRUE(e.matched);
@@ -264,7 +264,7 @@
TEST_F(WGSLParserTest, BitwiseExpr_And_InvalidRHS) {
auto p = parser("true & if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->bitwise_expression_post_unary_expression(lhs.value);
+ auto e = p->bitwise_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.matched);
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/expression_test.cc b/src/tint/lang/wgsl/reader/parser/expression_test.cc
index 2b6c7a2..74d78df 100644
--- a/src/tint/lang/wgsl/reader/parser/expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/expression_test.cc
@@ -50,9 +50,9 @@
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -117,9 +117,9 @@
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -180,7 +180,11 @@
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:3: mixing '&&' and '||' requires parenthesis");
+ EXPECT_EQ(p->builder().Diagnostics().str(),
+ R"(test.wgsl:1:3 error: mixing '&&' and '||' requires parenthesis
+a && true || b
+ ^^^^^^^^^^
+)");
}
TEST_F(WGSLParserTest, Expression_Mixing_AndWithOr) {
@@ -190,7 +194,11 @@
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:3: mixing '||' and '&&' requires parenthesis");
+ EXPECT_EQ(p->builder().Diagnostics().str(),
+ R"(test.wgsl:1:3 error: mixing '||' and '&&' requires parenthesis
+a || true && b
+ ^^^^^^^^^^
+)");
}
TEST_F(WGSLParserTest, Expression_Bitwise) {
@@ -281,7 +289,7 @@
EXPECT_TRUE(e.errored);
EXPECT_TRUE(p->has_error());
EXPECT_EQ(e.value, nullptr);
- EXPECT_EQ(p->error(), R"(1:7: mixing '&&' and '||' requires parenthesis)");
+ EXPECT_EQ(p->error(), R"(1:3: mixing '&&' and '||' requires parenthesis)");
}
TEST_F(WGSLParserTest, Expression_SubtractionNoSpace) {
diff --git a/src/tint/lang/wgsl/reader/parser/math_expression_test.cc b/src/tint/lang/wgsl/reader/parser/math_expression_test.cc
index 04e1165..4ae6f8e 100644
--- a/src/tint/lang/wgsl/reader/parser/math_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/math_expression_test.cc
@@ -33,7 +33,7 @@
TEST_F(WGSLParserTest, MathExpression_Parses_Multiplicative) {
auto p = parser("a * b");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -54,7 +54,7 @@
TEST_F(WGSLParserTest, MathExpression_Parses_Mixed_MultiplicativeStart) {
auto p = parser("a * b + c");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -89,7 +89,7 @@
TEST_F(WGSLParserTest, MathExpression_Parses_Additive) {
auto p = parser("a + b");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -110,7 +110,7 @@
TEST_F(WGSLParserTest, MathExpression_Parses_Mixed_AdditiveStart) {
auto p = parser("a + b * c");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -145,7 +145,7 @@
TEST_F(WGSLParserTest, MathExpression_NoMatch_ReturnLHS) {
auto p = parser("a if");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -155,7 +155,7 @@
TEST_F(WGSLParserTest, MathExpression_InvalidRHS) {
auto p = parser("a * if");
auto lhs = p->unary_expression();
- auto e = p->expect_math_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_math_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.errored);
EXPECT_TRUE(p->has_error());
ASSERT_EQ(e.value, nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/multiplicative_expression_test.cc b/src/tint/lang/wgsl/reader/parser/multiplicative_expression_test.cc
index b41347c..d0e11b7 100644
--- a/src/tint/lang/wgsl/reader/parser/multiplicative_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/multiplicative_expression_test.cc
@@ -33,7 +33,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_Parses_Multiply) {
auto p = parser("a * b");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -54,7 +54,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_Parses_Multiply_UnaryIndirect) {
auto p = parser("a **b");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -79,7 +79,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_Parses_Divide) {
auto p = parser("a / b");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -100,7 +100,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_Parses_Modulo) {
auto p = parser("a % b");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -121,7 +121,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_Parses_Grouping) {
auto p = parser("a * b / c % d * e");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -178,7 +178,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_InvalidRHS) {
auto p = parser("a * if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.errored);
EXPECT_EQ(e.value, nullptr);
ASSERT_TRUE(p->has_error());
@@ -188,7 +188,7 @@
TEST_F(WGSLParserTest, MultiplicativeExpression_NoMatch_ReturnsLHS) {
auto p = parser("a + b");
auto lhs = p->unary_expression();
- auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_multiplicative_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/parser.cc b/src/tint/lang/wgsl/reader/parser/parser.cc
index e3698f9..104f37d 100644
--- a/src/tint/lang/wgsl/reader/parser/parser.cc
+++ b/src/tint/lang/wgsl/reader/parser/parser.cc
@@ -173,7 +173,11 @@
/// @returns the Source that returns the combined source from start to the current last token's
/// source.
- tint::Source Source() const {
+ tint::Source operator()() const { return *this; }
+
+ /// @returns the Source that returns the combined source from start to the current last token's
+ /// source.
+ operator tint::Source() const {
auto end = parser_->last_source().End();
if (end < start_) {
end = start_;
@@ -181,10 +185,6 @@
return Source::Combine(start_, end);
}
- /// Implicit conversion to Source that returns the combined source from start to the current
- /// last token's source.
- operator tint::Source() const { return Source(); }
-
private:
Parser* parser_;
tint::Source start_;
@@ -447,7 +447,7 @@
return Failure::kErrored;
}
- builder_.AST().AddEnable(create<ast::Enable>(decl_source.Source(), std::move(extensions)));
+ builder_.AST().AddEnable(create<ast::Enable>(decl_source(), std::move(extensions)));
return kSuccess;
});
}
@@ -508,8 +508,7 @@
return Failure::kErrored;
}
- builder_.AST().AddRequires(
- create<ast::Requires>(decl_source.Source(), std::move(features)));
+ builder_.AST().AddRequires(create<ast::Requires>(decl_source(), std::move(features)));
return kSuccess;
});
}
@@ -857,8 +856,8 @@
// type_alias_decl
// : ALIAS IDENT EQUAL type_specifier
Maybe<const ast::Alias*> Parser::type_alias_decl() {
- Source source;
- if (!match(Token::Type::kAlias, &source)) {
+ MultiTokenSource source(this);
+ if (!match(Token::Type::kAlias)) {
return Failure::kNoMatch;
}
@@ -881,7 +880,7 @@
return add_error(peek(), "invalid type alias");
}
- return builder_.ty.alias(make_source_range_from(source), name.value, type.value);
+ return builder_.ty.alias(source(), name.value, type.value);
}
// type_specifier
@@ -894,7 +893,7 @@
}
if (!peek_is(Token::Type::kTemplateArgsLeft)) {
- return builder_.ty(builder_.Ident(source.Source(), ident.to_str()));
+ return builder_.ty(builder_.Ident(source(), ident.to_str()));
}
auto args = expect_template_arg_block("type template arguments", [&] {
@@ -904,7 +903,7 @@
if (args.errored) {
return Failure::kErrored;
}
- return builder_.ty(builder_.Ident(source.Source(), ident.to_str(), std::move(args.value)));
+ return builder_.ty(builder_.Ident(source(), ident.to_str(), std::move(args.value)));
}
template <typename ENUM>
@@ -1042,8 +1041,8 @@
// const_assert_statement
// : STATIC_ASSERT expression
Maybe<const ast::ConstAssert*> Parser::const_assert_statement() {
- Source start;
- if (!match(Token::Type::kConstAssert, &start)) {
+ MultiTokenSource source(this);
+ if (!match(Token::Type::kConstAssert)) {
return Failure::kNoMatch;
}
@@ -1055,8 +1054,7 @@
return add_error(peek(), "unable to parse condition expression");
}
- Source source = make_source_range_from(start);
- return create<ast::ConstAssert>(source, condition.value);
+ return create<ast::ConstAssert>(source(), condition.value);
}
// function_decl
@@ -1213,16 +1211,14 @@
// : attribute* BRACE_LEFT statement* BRACE_RIGHT
Expect<ast::BlockStatement*> Parser::expect_compound_statement(AttributeList& attrs,
std::string_view use) {
- auto source_start = peek().source();
+ MultiTokenSource source(this);
auto stmts =
expect_brace_block(use, [&]() -> Expect<StatementList> { return expect_statements(); });
- auto source_end = last_source();
if (stmts.errored) {
return Failure::kErrored;
}
TINT_DEFER(attrs.Clear());
- return create<ast::BlockStatement>(Source::Combine(source_start, source_end), stmts.value,
- std::move(attrs));
+ return create<ast::BlockStatement>(source(), stmts.value, std::move(attrs));
}
// paren_expression
@@ -1454,14 +1450,14 @@
// | LET optionally_typed_ident EQUAL expression
// | CONST optionally_typed_ident EQUAL expression
Maybe<const ast::VariableDeclStatement*> Parser::variable_statement() {
- auto decl_source_range = make_source_range();
+ MultiTokenSource decl_source_range(this);
if (match(Token::Type::kConst)) {
auto typed_ident = expect_optionally_typed_ident("'const' declaration");
if (typed_ident.errored) {
return Failure::kErrored;
}
- auto decl_source = decl_source_range.Source();
+ auto decl_source = decl_source_range();
if (!expect("'const' declaration", Token::Type::kEqual)) {
return Failure::kErrored;
@@ -1489,7 +1485,7 @@
return Failure::kErrored;
}
- auto decl_source = decl_source_range.Source();
+ auto decl_source = decl_source_range();
if (!expect("'let' declaration", Token::Type::kEqual)) {
return Failure::kErrored;
@@ -1519,7 +1515,7 @@
return Failure::kNoMatch;
}
- auto decl_source = decl_source_range.Source();
+ auto decl_source = decl_source_range();
const ast::Expression* initializer = nullptr;
if (match(Token::Type::kEqual)) {
@@ -1779,7 +1775,7 @@
}
Maybe<const ast::BlockStatement*> continuing(Failure::kErrored);
- auto body_start = peek().source();
+ MultiTokenSource body_source(this);
auto body = expect_brace_block("loop", [&]() -> Maybe<StatementList> {
auto stmts = expect_statements();
if (stmts.errored) {
@@ -1795,13 +1791,10 @@
if (body.errored) {
return Failure::kErrored;
}
- auto body_end = last_source();
TINT_DEFER(attrs.Clear());
return create<ast::LoopStatement>(
- source,
- create<ast::BlockStatement>(Source::Combine(body_start, body_end), body.value,
- std::move(body_attrs.value)),
+ source, create<ast::BlockStatement>(body_source(), body.value, std::move(body_attrs.value)),
continuing.value, std::move(attrs));
}
@@ -2017,7 +2010,7 @@
return Failure::kErrored;
}
- auto source_start = peek().source();
+ MultiTokenSource source(this);
auto body = expect_brace_block("", [&]() -> Expect<StatementList> {
StatementList stmts;
@@ -2047,10 +2040,8 @@
if (body.errored) {
return Failure::kErrored;
}
- auto source_end = last_source();
- return create<ast::BlockStatement>(Source::Combine(source_start, source_end), body.value,
- std::move(attrs.value));
+ return create<ast::BlockStatement>(source(), body.value, std::move(attrs.value));
}
// continuing_statement
@@ -2109,9 +2100,9 @@
return expect_expression_list("template argument list",
Token::Type::kTemplateArgsRight);
});
- ident = builder_.Ident(source.Source(), t.to_str(), std::move(tmpl_args.value));
+ ident = builder_.Ident(source(), t.to_str(), std::move(tmpl_args.value));
} else {
- ident = builder_.Ident(source.Source(), t.to_str());
+ ident = builder_.Ident(source(), t.to_str());
}
if (peek_is(Token::Type::kParenLeft)) {
@@ -2120,7 +2111,7 @@
return Failure::kErrored;
}
- return builder_.Call(source.Source(), ident, std::move(params.value));
+ return builder_.Call(source(), ident, std::move(params.value));
}
return builder_.Expr(ident);
@@ -2162,7 +2153,7 @@
return Failure::kErrored;
}
- return create<ast::IndexAccessorExpression>(source.Source(), prefix, param.value);
+ return create<ast::IndexAccessorExpression>(source(), prefix, param.value);
});
if (res.errored) {
@@ -2178,7 +2169,7 @@
return Failure::kErrored;
}
- prefix = builder_.MemberAccessor(source.Source(), prefix, ident.value);
+ prefix = builder_.MemberAccessor(source(), prefix, ident.value);
continue;
}
@@ -2205,7 +2196,9 @@
// | OR unary_expression (OR unary_expression)*
// | XOR unary_expression (XOR unary_expression)*
Maybe<const ast::Expression*> Parser::bitwise_expression_post_unary_expression(
- const ast::Expression* lhs) {
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
auto& t = peek();
std::optional<core::BinaryOp> op;
@@ -2234,7 +2227,7 @@
std::string(t.to_name()) + " expression");
}
- lhs = create<ast::BinaryExpression>(t.source(), *op, lhs, rhs.value);
+ lhs = create<ast::BinaryExpression>(source(), *op, lhs, rhs.value);
if (!match(t.type())) {
return lhs;
@@ -2264,7 +2257,9 @@
// multiplicative_expression.post.unary_expression
// : (multiplicative_operator unary_expression)*
Expect<const ast::Expression*> Parser::expect_multiplicative_expression_post_unary_expression(
- const ast::Expression* lhs) {
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
while (continue_parsing()) {
auto& t = peek();
@@ -2285,7 +2280,7 @@
std::string(t.to_name()) + " expression");
}
- lhs = create<ast::BinaryExpression>(t.source(), op.value, lhs, rhs.value);
+ lhs = create<ast::BinaryExpression>(source(), op.value, lhs, rhs.value);
}
return Failure::kErrored;
}
@@ -2320,7 +2315,9 @@
// This is `( additive_operator unary_expression ( multiplicative_operator unary_expression )* )*`
// split apart.
Expect<const ast::Expression*> Parser::expect_additive_expression_post_unary_expression(
- const ast::Expression* lhs) {
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
while (continue_parsing()) {
auto& t = peek();
@@ -2341,14 +2338,14 @@
std::string(t.to_name()) + " expression");
}
- // The multiplicative binds tigher, so pass the unary into that and build that expression
- // before creating the additve expression.
- auto rhs = expect_multiplicative_expression_post_unary_expression(unary.value);
+ // The multiplicative binds tighter, so pass the unary into that and build that expression
+ // before creating the additive expression.
+ auto rhs = expect_multiplicative_expression_post_unary_expression(unary.value, lhs_source);
if (rhs.errored) {
return Failure::kErrored;
}
- lhs = create<ast::BinaryExpression>(t.source(), op.value, lhs, rhs.value);
+ lhs = create<ast::BinaryExpression>(source(), op.value, lhs, rhs.value);
}
return Failure::kErrored;
}
@@ -2359,18 +2356,22 @@
// This is `( multiplicative_operator unary_expression )* ( additive_operator unary_expression (
// multiplicative_operator unary_expression )* )*` split apart.
Expect<const ast::Expression*> Parser::expect_math_expression_post_unary_expression(
- const ast::Expression* lhs) {
- auto rhs = expect_multiplicative_expression_post_unary_expression(lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
+
+ auto rhs = expect_multiplicative_expression_post_unary_expression(lhs, source);
if (rhs.errored) {
return Failure::kErrored;
}
- return expect_additive_expression_post_unary_expression(rhs.value);
+ return expect_additive_expression_post_unary_expression(rhs.value, source());
}
// shift_expression
// : unary_expression shift_expression.post.unary_expression
Maybe<const ast::Expression*> Parser::shift_expression() {
+ MultiTokenSource source(this);
auto lhs = unary_expression();
if (lhs.errored) {
return Failure::kErrored;
@@ -2378,7 +2379,7 @@
if (!lhs.matched) {
return Failure::kNoMatch;
}
- return expect_shift_expression_post_unary_expression(lhs.value);
+ return expect_shift_expression_post_unary_expression(lhs.value, source);
}
// shift_expression.post.unary_expression
@@ -2389,7 +2390,10 @@
// Note, add the `math_expression.post.unary_expression` is added here to make
// implementation simpler.
Expect<const ast::Expression*> Parser::expect_shift_expression_post_unary_expression(
- const ast::Expression* lhs) {
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
+
auto& t = peek();
if (match(Token::Type::kShiftLeft) || match(Token::Type::kShiftRight)) {
std::string name;
@@ -2411,15 +2415,16 @@
return add_error(rhs_start,
std::string("unable to parse right side of ") + name + " expression");
}
- return create<ast::BinaryExpression>(t.source(), *op, lhs, rhs.value);
+ return create<ast::BinaryExpression>(source(), *op, lhs, rhs.value);
}
- return expect_math_expression_post_unary_expression(lhs);
+ return expect_math_expression_post_unary_expression(lhs, source);
}
// relational_expression
// : unary_expression relational_expression.post.unary_expression
Maybe<const ast::Expression*> Parser::relational_expression() {
+ MultiTokenSource source(this);
auto lhs = unary_expression();
if (lhs.errored) {
return Failure::kErrored;
@@ -2427,7 +2432,7 @@
if (!lhs.matched) {
return Failure::kNoMatch;
}
- return expect_relational_expression_post_unary_expression(lhs.value);
+ return expect_relational_expression_post_unary_expression(lhs.value, source);
}
// relational_expression.post.unary_expression
@@ -2441,8 +2446,11 @@
//
// Note, a `shift_expression` element was added to simplify many of the right sides
Expect<const ast::Expression*> Parser::expect_relational_expression_post_unary_expression(
- const ast::Expression* lhs) {
- auto lhs_result = expect_shift_expression_post_unary_expression(lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source) {
+ MultiTokenSource source(this, lhs_source);
+
+ auto lhs_result = expect_shift_expression_post_unary_expression(lhs, source);
if (lhs_result.errored) {
return Failure::kErrored;
}
@@ -2486,7 +2494,7 @@
std::string(tok_op.to_name()) + " expression");
}
- return create<ast::BinaryExpression>(tok_op.source(), *op, lhs, rhs.value);
+ return create<ast::BinaryExpression>(source(), *op, lhs, rhs.value);
}
Expect<const ast::Expression*> Parser::expect_expression(std::string_view use) {
@@ -2560,6 +2568,8 @@
//
// Note, a `relational_expression` element was added to simplify many of the right sides
Maybe<const ast::Expression*> Parser::expression() {
+ MultiTokenSource source(this);
+ Source first_op;
auto expr = [&]() -> Maybe<const ast::Expression*> {
auto lhs = unary_expression();
if (lhs.errored) {
@@ -2569,7 +2579,9 @@
return Failure::kNoMatch;
}
- auto bitwise = bitwise_expression_post_unary_expression(lhs.value);
+ first_op = peek().source();
+
+ auto bitwise = bitwise_expression_post_unary_expression(lhs.value, source);
if (bitwise.errored) {
return Failure::kErrored;
}
@@ -2577,7 +2589,7 @@
return bitwise.value;
}
- auto relational = expect_relational_expression_post_unary_expression(lhs.value);
+ auto relational = expect_relational_expression_post_unary_expression(lhs.value, source);
if (relational.errored) {
return Failure::kErrored;
}
@@ -2608,7 +2620,7 @@
std::string(t.to_name()) + " expression");
}
- ret = create<ast::BinaryExpression>(t.source(), op, ret, rhs.value);
+ ret = create<ast::BinaryExpression>(source(), op, ret, rhs.value);
}
}
return ret;
@@ -2620,9 +2632,9 @@
// after this then it _must_ be a different one, and hence an error.
if (auto* lhs = expr->As<ast::BinaryExpression>()) {
if (auto& n = peek(); n.IsBinaryOperator()) {
- auto source = Source::Combine(expr->source, n.source());
- add_error(source, std::string("mixing '") + ast::Operator(lhs->op) + "' and '" +
- std::string(n.to_name()) + "' requires parenthesis");
+ add_error(Source::Combine(first_op, n.source()),
+ std::string("mixing '") + ast::Operator(lhs->op) + "' and '" +
+ std::string(n.to_name()) + "' requires parenthesis");
return Failure::kErrored;
}
}
@@ -2656,10 +2668,11 @@
// The `primary_expression component_or_swizzle_specifier ?` is moved out into a
// `singular_expression`
Maybe<const ast::Expression*> Parser::unary_expression() {
- auto& t = peek();
+ MultiTokenSource source(this);
+ auto& t = peek();
if (match(Token::Type::kPlusPlus) || match(Token::Type::kMinusMinus)) {
- add_error(t.source(),
+ add_error(source,
"prefix increment and decrement operators are reserved for a "
"future WGSL version");
return Failure::kErrored;
@@ -2700,7 +2713,7 @@
peek(), "unable to parse right side of " + std::string(t.to_name()) + " expression");
}
- return create<ast::UnaryOpExpression>(t.source(), op, expr.value);
+ return create<ast::UnaryOpExpression>(source(), op, expr.value);
}
// compound_assignment_operator
@@ -3500,12 +3513,4 @@
return result;
}
-Parser::MultiTokenSource Parser::make_source_range() {
- return MultiTokenSource(this);
-}
-
-Parser::MultiTokenSource Parser::make_source_range_from(const Source& start) {
- return MultiTokenSource(this, start);
-}
-
} // namespace tint::wgsl::reader
diff --git a/src/tint/lang/wgsl/reader/parser/parser.h b/src/tint/lang/wgsl/reader/parser/parser.h
index 0e6c152..58f231e 100644
--- a/src/tint/lang/wgsl/reader/parser/parser.h
+++ b/src/tint/lang/wgsl/reader/parser/parser.h
@@ -586,43 +586,55 @@
Token::Type terminator);
/// Parses the `bitwise_expression.post.unary_expression` grammar element
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or nullptr
Maybe<const ast::Expression*> bitwise_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parse the `multiplicative_operator` grammar element
/// @returns the parsed operator if successful
Maybe<core::BinaryOp> multiplicative_operator();
/// Parses multiplicative elements
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or `lhs` if no match
Expect<const ast::Expression*> expect_multiplicative_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parses additive elements
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or `lhs` if no match
Expect<const ast::Expression*> expect_additive_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parses math elements
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or `lhs` if no match
Expect<const ast::Expression*> expect_math_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parses a `unary_expression shift.post.unary_expression`
/// @returns the parsed expression or nullptr
Maybe<const ast::Expression*> shift_expression();
/// Parses a `shift_expression.post.unary_expression` grammar element
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or `lhs` if no match
Expect<const ast::Expression*> expect_shift_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parses a `unary_expression relational_expression.post.unary_expression`
/// @returns the parsed expression or nullptr
Maybe<const ast::Expression*> relational_expression();
/// Parses a `relational_expression.post.unary_expression` grammar element
/// @param lhs the left side of the expression
+ /// @param lhs_source the source span for the left side of the expression
/// @returns the parsed expression or `lhs` if no match
Expect<const ast::Expression*> expect_relational_expression_post_unary_expression(
- const ast::Expression* lhs);
+ const ast::Expression* lhs,
+ const Source& lhs_source);
/// Parse the `additive_operator` grammar element
/// @returns the parsed operator if successful
Maybe<core::BinaryOp> additive_operator();
@@ -881,8 +893,6 @@
Maybe<const ast::Statement*> for_header_continuing();
class MultiTokenSource;
- MultiTokenSource make_source_range();
- MultiTokenSource make_source_range_from(const Source& start);
/// Creates a new `ast::Node` owned by the Module. When the Module is
/// destructed, the `ast::Node` will also be destructed.
diff --git a/src/tint/lang/wgsl/reader/parser/relational_expression_test.cc b/src/tint/lang/wgsl/reader/parser/relational_expression_test.cc
index e2122ca..97b0d91 100644
--- a/src/tint/lang/wgsl/reader/parser/relational_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/relational_expression_test.cc
@@ -33,15 +33,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_LessThan) {
auto p = parser("a < true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 9u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -58,15 +58,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_GreaterThan) {
auto p = parser("a > true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 4u);
+ EXPECT_EQ(e->source.range.end.column, 9u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -83,15 +83,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_LessThanEqual) {
auto p = parser("a <= true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -108,15 +108,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_GreaterThanEqual) {
auto p = parser("a >= true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -133,15 +133,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_Equal) {
auto p = parser("a == true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -158,15 +158,15 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_Parses_NotEqual) {
auto p = parser("a != true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -183,7 +183,7 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_InvalidRHS) {
auto p = parser("true < if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(e.value, nullptr);
EXPECT_EQ(p->error(), "1:8: unable to parse right side of < expression");
@@ -192,7 +192,7 @@
TEST_F(WGSLParserTest, RelationalExpression_PostUnary_NoMatch_ReturnsLHS) {
auto p = parser("a true");
auto lhs = p->unary_expression();
- auto e = p->expect_relational_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_relational_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -247,9 +247,9 @@
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -272,9 +272,9 @@
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
diff --git a/src/tint/lang/wgsl/reader/parser/shift_expression_test.cc b/src/tint/lang/wgsl/reader/parser/shift_expression_test.cc
index 77ec7db..011b131 100644
--- a/src/tint/lang/wgsl/reader/parser/shift_expression_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/shift_expression_test.cc
@@ -33,15 +33,15 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_Parses_ShiftLeft) {
auto p = parser("a << true");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -58,15 +58,15 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_Parses_ShiftRight) {
auto p = parser("a >> true");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
EXPECT_EQ(e->source.range.begin.line, 1u);
- EXPECT_EQ(e->source.range.begin.column, 3u);
+ EXPECT_EQ(e->source.range.begin.column, 1u);
EXPECT_EQ(e->source.range.end.line, 1u);
- EXPECT_EQ(e->source.range.end.column, 5u);
+ EXPECT_EQ(e->source.range.end.column, 10u);
ASSERT_TRUE(e->Is<ast::BinaryExpression>());
auto* rel = e->As<ast::BinaryExpression>();
@@ -83,7 +83,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_Parses_Additive) {
auto p = parser("a + b");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -104,7 +104,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_Parses_Multiplicative) {
auto p = parser("a * b");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
@@ -125,7 +125,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_InvalidSpaceLeft) {
auto p = parser("a < < true");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
ASSERT_NE(e.value, nullptr);
EXPECT_FALSE(e.value->Is<ast::BinaryExpression>());
@@ -134,7 +134,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_InvalidSpaceRight) {
auto p = parser("a > > true");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
ASSERT_NE(e.value, nullptr);
EXPECT_FALSE(e.value->Is<ast::BinaryExpression>());
@@ -143,7 +143,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_InvalidRHS) {
auto p = parser("a << if (a) {}");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_TRUE(e.errored);
EXPECT_TRUE(p->has_error());
EXPECT_EQ(e.value, nullptr);
@@ -153,7 +153,7 @@
TEST_F(WGSLParserTest, ShiftExpression_PostUnary_NoOr_ReturnsLHS) {
auto p = parser("a true");
auto lhs = p->unary_expression();
- auto e = p->expect_shift_expression_post_unary_expression(lhs.value);
+ auto e = p->expect_shift_expression_post_unary_expression(lhs.value, lhs->source);
EXPECT_FALSE(e.errored);
EXPECT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e.value, nullptr);
diff --git a/src/tint/lang/wgsl/reader/parser/struct_member_attribute_test.cc b/src/tint/lang/wgsl/reader/parser/struct_member_attribute_test.cc
index 50b7908..f4169ac 100644
--- a/src/tint/lang/wgsl/reader/parser/struct_member_attribute_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/struct_member_attribute_test.cc
@@ -233,7 +233,11 @@
EXPECT_EQ(attr.value, nullptr);
EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:9: mixing '+' and '<<' requires parenthesis");
+ EXPECT_EQ(p->builder().Diagnostics().str(),
+ R"(test.wgsl:1:9 error: mixing '+' and '<<' requires parenthesis
+align(4 + 5 << 6)
+ ^^^^^^
+)");
}
TEST_F(WGSLParserTest, Attribute_Index) {
diff --git a/src/tint/lang/wgsl/resolver/uniformity.cc b/src/tint/lang/wgsl/resolver/uniformity.cc
index b3dc470..709da02 100644
--- a/src/tint/lang/wgsl/resolver/uniformity.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity.cc
@@ -1232,10 +1232,18 @@
node->AddEdge(cf);
auto* current_value = current_function_->variables.Get(param);
- if (param->Type()->Is<core::type::Pointer>()) {
+ if (auto* ptr = param->Type()->As<core::type::Pointer>()) {
if (load_rule) {
- // We are loading from the pointer, so add an edge to its contents.
- node->AddEdge(current_value);
+ if (ptr->AddressSpace() == core::AddressSpace::kFunction ||
+ ptr->Access() == core::Access::kRead) {
+ // We are loading from a pointer to a function-scope variable or an
+ // immutable module-scope variable, so add an edge to its contents.
+ node->AddEdge(current_value);
+ } else {
+ // We are loading from a pointer to a mutable module-scope variable,
+ // which always has non-uniform contents.
+ node->AddEdge(current_function_->may_be_non_uniform);
+ }
} else {
// This is a pointer parameter that we are not loading from, so add an
// edge to the pointer value itself.
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index 422c711..cbab7db 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -2543,9 +2543,9 @@
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:7:34 note: control flow depends on possibly non-uniform value
+test:7:7 note: control flow depends on possibly non-uniform value
if ((non_uniform_global == 42) && false) {
- ^^
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
test:7:8 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
if ((non_uniform_global == 42) && false) {
@@ -2601,9 +2601,9 @@
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:7:34 note: control flow depends on possibly non-uniform value
+test:7:7 note: control flow depends on possibly non-uniform value
if ((non_uniform_global == 42) || true) {
- ^^
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
test:7:8 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
if ((non_uniform_global == 42) || true) {
@@ -3750,7 +3750,7 @@
test:12:7 note: possibly non-uniform value passed via pointer here
bar(&v);
- ^
+ ^^
test:11:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
var v = non_uniform;
@@ -3816,14 +3816,6 @@
test:7:8 note: parameter 'p' of 'bar' may be non-uniform
if (*p == 0) {
^
-
-test:13:7 note: possibly non-uniform value passed via pointer here
- bar(&non_uniform);
- ^
-
-test:4:48 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
-@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
- ^^^^^^^^^^^
)");
}
@@ -3854,9 +3846,9 @@
if (0 == bar(&non_uniform)) {
^^
-test:4:48 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
-@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
- ^^^^^^^^^^^
+test:11:12 note: return value of 'bar' may be non-uniform
+ if (0 == bar(&non_uniform)) {
+ ^^^^^^^^^^^^^^^^^
)");
}
@@ -3893,7 +3885,7 @@
test:13:7 note: possibly non-uniform value passed via pointer here
bar(&v);
- ^
+ ^^
test:12:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
var v = non_uniform;
@@ -3934,7 +3926,7 @@
test:12:7 note: possibly non-uniform value passed via pointer here
bar(&v);
- ^
+ ^^
test:11:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
var v = non_uniform;
@@ -4439,7 +4431,7 @@
test:10:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -4703,7 +4695,7 @@
test:16:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -4740,7 +4732,7 @@
test:14:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -4777,7 +4769,7 @@
test:14:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -4822,7 +4814,7 @@
test:22:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -4861,7 +4853,7 @@
test:16:7 note: contents of pointer may become non-uniform after calling 'bar'
bar(&v);
- ^
+ ^^
)");
}
@@ -5164,7 +5156,7 @@
test:12:11 note: contents of pointer may become non-uniform after calling 'bar'
bar(&a, &b);
- ^
+ ^^
)");
}
@@ -5257,7 +5249,7 @@
test:12:11 note: contents of pointer may become non-uniform after calling 'bar'
bar(&a, &b);
- ^
+ ^^
)");
}
@@ -5361,6 +5353,128 @@
note: reading from module-scope private variable 'non_uniform_global' may result in a non-uniform value)");
}
+TEST_F(UniformityAnalysisTest, AssignUniformToPrivatePointerParameter_StillNonUniform) {
+ std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> non_uniform : i32;
+
+fn bar(p : ptr<private, i32>) {
+ *p = 0;
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn foo() {
+ bar(&non_uniform);
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:8:3 note: control flow depends on possibly non-uniform value
+ if (*p == 0) {
+ ^^
+
+test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+ if (*p == 0) {
+ ^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformToWorkgroupPointerParameter_StillNonUniform) {
+ std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> non_uniform : i32;
+
+fn bar(p : ptr<workgroup, i32>) {
+ *p = 0;
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn foo() {
+ bar(&non_uniform);
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:8:3 note: control flow depends on possibly non-uniform value
+ if (*p == 0) {
+ ^^
+
+test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+ if (*p == 0) {
+ ^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformToStoragePointerParameter_StillNonUniform) {
+ std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<storage, i32, read_write>) {
+ *p = 0;
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn foo() {
+ bar(&non_uniform);
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:8:3 note: control flow depends on possibly non-uniform value
+ if (*p == 0) {
+ ^^
+
+test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+ if (*p == 0) {
+ ^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, LoadFromReadOnlyStoragePointerParameter_AlwaysUniform) {
+ std::string src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read> non_uniform : i32;
+
+fn bar(p : ptr<storage, i32, read>) {
+ if (*p == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn foo() {
+ bar(&non_uniform);
+}
+)";
+
+ RunTest(src, true);
+}
+
////////////////////////////////////////////////////////////////////////////////
/// Tests to cover access to aggregate types.
////////////////////////////////////////////////////////////////////////////////
@@ -7944,7 +8058,7 @@
test:15:9 note: possibly non-uniform value passed via pointer here
v[bar(&f)] += 1;
- ^
+ ^^
test:14:11 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
var f = rw;
@@ -8020,7 +8134,7 @@
test:15:9 note: possibly non-uniform value passed via pointer here
v[bar(&f)]++;
- ^
+ ^^
test:14:11 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
var f = rw;
@@ -8093,11 +8207,11 @@
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)];
- ^
+ ^^
)");
}
@@ -8330,11 +8444,11 @@
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)];
- ^
+ ^^
)");
}
@@ -8468,9 +8582,9 @@
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
^^^^^^^^^
-test:5:37 note: control flow depends on possibly non-uniform value
+test:5:11 note: control flow depends on possibly non-uniform value
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
- ^^
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
test:5:12 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
@@ -8725,9 +8839,9 @@
if (0 == bar(p)) {
^^
-test:4:48 note: reading from read_write storage buffer 'arr' may result in a non-uniform value
-@group(0) @binding(0) var<storage, read_write> arr : array<u32>;
- ^^^
+test:19:12 note: return value of 'bar' may be non-uniform
+ if (0 == bar(p)) {
+ ^^^^^^
)");
}
@@ -8789,7 +8903,7 @@
EXPECT_EQ(error_,
R"(test:8:28 error: possibly non-uniform value passed here
if (workgroupUniformLoad(&data[idx]) > 0) {
- ^
+ ^^^^^^^^^^
test:8:34 note: builtin 'idx' of 'main' may be non-uniform
if (workgroupUniformLoad(&data[idx]) > 0) {
@@ -8829,7 +8943,7 @@
test:14:11 note: possibly non-uniform value passed here
if (foo(&data[idx]) > 0) {
- ^
+ ^^^^^^^^^^
test:14:17 note: builtin 'idx' of 'main' may be non-uniform
if (foo(&data[idx]) > 0) {
