[tint][ast] Fix dynamic indices in DirectVariableAccess
The logic to extract the callee's dynamic indices was incorrect.
Add more tests, for both the AST and IR transforms.
Bug: tint:2053
Change-Id: I9fa2eeded5a9ba525b1529377a0ab443e80d8634
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/169263
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/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/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/test/tint/array/strides.spvasm.expected.glsl b/test/tint/array/strides.spvasm.expected.glsl
index b82b24f..94ee04c 100644
--- a/test/tint/array/strides.spvasm.expected.glsl
+++ b/test/tint/array/strides.spvasm.expected.glsl
@@ -38,7 +38,7 @@
} s;
void assign_and_preserve_padding_4_s_a_X_el_X_X(uint dest[3], strided_arr value) {
- s.inner.a[dest[0]].el[dest[0]][dest[0]].el = value.el;
+ s.inner.a[dest[0]].el[dest[1]][dest[2]].el = value.el;
}
void assign_and_preserve_padding_3_s_a_X_el_X(uint dest[2], strided_arr value[2]) {
diff --git a/test/tint/array/strides.spvasm.expected.spvasm b/test/tint/array/strides.spvasm.expected.spvasm
index 3d87847..d319985 100644
--- a/test/tint/array/strides.spvasm.expected.spvasm
+++ b/test/tint/array/strides.spvasm.expected.spvasm
@@ -71,181 +71,181 @@
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%24 = OpConstantNull %int
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
- %31 = OpTypeFunction %void %_arr_uint_uint_2 %_arr_strided_arr_uint_2
- %37 = OpConstantNull %uint
+ %33 = OpTypeFunction %void %_arr_uint_uint_2 %_arr_strided_arr_uint_2
+ %39 = OpConstantNull %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%bool = OpTypeBool
%_ptr_Function__arr_strided_arr_uint_2 = OpTypePointer Function %_arr_strided_arr_uint_2
- %52 = OpConstantNull %_arr_strided_arr_uint_2
+ %54 = OpConstantNull %_arr_strided_arr_uint_2
%uint_1 = OpConstant %uint 1
%_ptr_Function_strided_arr = OpTypePointer Function %strided_arr
%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
- %65 = OpTypeFunction %void %_arr_uint_uint_1 %_arr__arr_strided_arr_uint_2_uint_3
+ %67 = OpTypeFunction %void %_arr_uint_uint_1 %_arr__arr_strided_arr_uint_2_uint_3
%_ptr_Function__arr__arr_strided_arr_uint_2_uint_3 = OpTypePointer Function %_arr__arr_strided_arr_uint_2_uint_3
- %83 = OpConstantNull %_arr__arr_strided_arr_uint_2_uint_3
- %93 = OpTypeFunction %void %_arr_uint_uint_1 %strided_arr_1
- %102 = OpTypeFunction %void %_arr_strided_arr_1_uint_4
+ %85 = OpConstantNull %_arr__arr_strided_arr_uint_2_uint_3
+ %95 = OpTypeFunction %void %_arr_uint_uint_1 %strided_arr_1
+ %104 = OpTypeFunction %void %_arr_strided_arr_1_uint_4
%_ptr_Function__arr_strided_arr_1_uint_4 = OpTypePointer Function %_arr_strided_arr_1_uint_4
- %118 = OpConstantNull %_arr_strided_arr_1_uint_4
+ %120 = OpConstantNull %_arr_strided_arr_1_uint_4
%_ptr_Function_strided_arr_1 = OpTypePointer Function %strided_arr_1
- %128 = OpTypeFunction %void
+ %130 = OpTypeFunction %void
%_ptr_StorageBuffer__arr_strided_arr_1_uint_4 = OpTypePointer StorageBuffer %_arr_strided_arr_1_uint_4
%int_3 = OpConstant %int 3
%_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 = OpTypePointer StorageBuffer %_arr__arr_strided_arr_uint_2_uint_3
- %int_2 = OpConstant %int 2
%_ptr_StorageBuffer__arr_strided_arr_uint_2 = OpTypePointer StorageBuffer %_arr_strided_arr_uint_2
- %int_1 = OpConstant %int 1
%float_5 = OpConstant %float 5
%assign_and_preserve_padding_4_s_a_X_el_X_X = OpFunction %void None %15
%dest = OpFunctionParameter %_arr_uint_uint_3
%value = OpFunctionParameter %strided_arr
%21 = OpLabel
%25 = OpCompositeExtract %uint %dest 0
- %26 = OpCompositeExtract %uint %dest 0
- %27 = OpCompositeExtract %uint %dest 0
- %29 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %uint_0 %25 %uint_0 %26 %27 %uint_0
- %30 = OpCompositeExtract %float %value 0
- OpStore %29 %30
+ %27 = OpCompositeExtract %uint %dest 1
+ %29 = OpCompositeExtract %uint %dest 2
+ %31 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %uint_0 %25 %uint_0 %27 %29 %uint_0
+ %32 = OpCompositeExtract %float %value 0
+ OpStore %31 %32
OpReturn
OpFunctionEnd
-%assign_and_preserve_padding_3_s_a_X_el_X = OpFunction %void None %31
+%assign_and_preserve_padding_3_s_a_X_el_X = OpFunction %void None %33
%dest_0 = OpFunctionParameter %_arr_uint_uint_2
%value_0 = OpFunctionParameter %_arr_strided_arr_uint_2
- %36 = OpLabel
- %i = OpVariable %_ptr_Function_uint Function %37
-%var_for_index = OpVariable %_ptr_Function__arr_strided_arr_uint_2 Function %52
- OpStore %i %37
- OpBranch %40
- %40 = OpLabel
- OpLoopMerge %41 %42 None
- OpBranch %43
- %43 = OpLabel
- %45 = OpLoad %uint %i
- %46 = OpULessThan %bool %45 %uint_2
- %44 = OpLogicalNot %bool %46
- OpSelectionMerge %48 None
- OpBranchConditional %44 %49 %48
- %49 = OpLabel
- OpBranch %41
- %48 = OpLabel
- OpStore %var_for_index %value_0
- %54 = OpCompositeExtract %uint %dest_0 0
- %56 = OpCompositeExtract %uint %dest_0 1
- %57 = OpLoad %uint %i
- %58 = OpCompositeConstruct %_arr_uint_uint_3 %54 %56 %57
- %59 = OpLoad %uint %i
- %61 = OpAccessChain %_ptr_Function_strided_arr %var_for_index %59
- %62 = OpLoad %strided_arr %61
- %53 = OpFunctionCall %void %assign_and_preserve_padding_4_s_a_X_el_X_X %58 %62
+ %38 = OpLabel
+ %i = OpVariable %_ptr_Function_uint Function %39
+%var_for_index = OpVariable %_ptr_Function__arr_strided_arr_uint_2 Function %54
+ OpStore %i %39
OpBranch %42
%42 = OpLabel
- %63 = OpLoad %uint %i
- %64 = OpIAdd %uint %63 %uint_1
- OpStore %i %64
- OpBranch %40
- %41 = OpLabel
+ OpLoopMerge %43 %44 None
+ OpBranch %45
+ %45 = OpLabel
+ %47 = OpLoad %uint %i
+ %48 = OpULessThan %bool %47 %uint_2
+ %46 = OpLogicalNot %bool %48
+ OpSelectionMerge %50 None
+ OpBranchConditional %46 %51 %50
+ %51 = OpLabel
+ OpBranch %43
+ %50 = OpLabel
+ OpStore %var_for_index %value_0
+ %56 = OpCompositeExtract %uint %dest_0 0
+ %58 = OpCompositeExtract %uint %dest_0 1
+ %59 = OpLoad %uint %i
+ %60 = OpCompositeConstruct %_arr_uint_uint_3 %56 %58 %59
+ %61 = OpLoad %uint %i
+ %63 = OpAccessChain %_ptr_Function_strided_arr %var_for_index %61
+ %64 = OpLoad %strided_arr %63
+ %55 = OpFunctionCall %void %assign_and_preserve_padding_4_s_a_X_el_X_X %60 %64
+ OpBranch %44
+ %44 = OpLabel
+ %65 = OpLoad %uint %i
+ %66 = OpIAdd %uint %65 %uint_1
+ OpStore %i %66
+ OpBranch %42
+ %43 = OpLabel
OpReturn
OpFunctionEnd
-%assign_and_preserve_padding_2_s_a_X_el = OpFunction %void None %65
+%assign_and_preserve_padding_2_s_a_X_el = OpFunction %void None %67
%dest_1 = OpFunctionParameter %_arr_uint_uint_1
%value_1 = OpFunctionParameter %_arr__arr_strided_arr_uint_2_uint_3
- %70 = OpLabel
- %i_0 = OpVariable %_ptr_Function_uint Function %37
-%var_for_index_1 = OpVariable %_ptr_Function__arr__arr_strided_arr_uint_2_uint_3 Function %83
- OpStore %i_0 %37
- OpBranch %72
%72 = OpLabel
- OpLoopMerge %73 %74 None
- OpBranch %75
- %75 = OpLabel
- %77 = OpLoad %uint %i_0
- %78 = OpULessThan %bool %77 %uint_3
- %76 = OpLogicalNot %bool %78
- OpSelectionMerge %79 None
- OpBranchConditional %76 %80 %79
- %80 = OpLabel
- OpBranch %73
- %79 = OpLabel
- OpStore %var_for_index_1 %value_1
- %85 = OpCompositeExtract %uint %dest_1 0
- %86 = OpLoad %uint %i_0
- %87 = OpCompositeConstruct %_arr_uint_uint_2 %85 %86
- %88 = OpLoad %uint %i_0
- %89 = OpAccessChain %_ptr_Function__arr_strided_arr_uint_2 %var_for_index_1 %88
- %90 = OpLoad %_arr_strided_arr_uint_2 %89
- %84 = OpFunctionCall %void %assign_and_preserve_padding_3_s_a_X_el_X %87 %90
+ %i_0 = OpVariable %_ptr_Function_uint Function %39
+%var_for_index_1 = OpVariable %_ptr_Function__arr__arr_strided_arr_uint_2_uint_3 Function %85
+ OpStore %i_0 %39
OpBranch %74
%74 = OpLabel
- %91 = OpLoad %uint %i_0
- %92 = OpIAdd %uint %91 %uint_1
- OpStore %i_0 %92
- OpBranch %72
- %73 = OpLabel
+ OpLoopMerge %75 %76 None
+ OpBranch %77
+ %77 = OpLabel
+ %79 = OpLoad %uint %i_0
+ %80 = OpULessThan %bool %79 %uint_3
+ %78 = OpLogicalNot %bool %80
+ OpSelectionMerge %81 None
+ OpBranchConditional %78 %82 %81
+ %82 = OpLabel
+ OpBranch %75
+ %81 = OpLabel
+ OpStore %var_for_index_1 %value_1
+ %87 = OpCompositeExtract %uint %dest_1 0
+ %88 = OpLoad %uint %i_0
+ %89 = OpCompositeConstruct %_arr_uint_uint_2 %87 %88
+ %90 = OpLoad %uint %i_0
+ %91 = OpAccessChain %_ptr_Function__arr_strided_arr_uint_2 %var_for_index_1 %90
+ %92 = OpLoad %_arr_strided_arr_uint_2 %91
+ %86 = OpFunctionCall %void %assign_and_preserve_padding_3_s_a_X_el_X %89 %92
+ OpBranch %76
+ %76 = OpLabel
+ %93 = OpLoad %uint %i_0
+ %94 = OpIAdd %uint %93 %uint_1
+ OpStore %i_0 %94
+ OpBranch %74
+ %75 = OpLabel
OpReturn
OpFunctionEnd
-%assign_and_preserve_padding_1_s_a_X = OpFunction %void None %93
+%assign_and_preserve_padding_1_s_a_X = OpFunction %void None %95
%dest_2 = OpFunctionParameter %_arr_uint_uint_1
%value_2 = OpFunctionParameter %strided_arr_1
- %97 = OpLabel
- %99 = OpCompositeExtract %uint %dest_2 0
- %100 = OpCompositeConstruct %_arr_uint_uint_1 %99
- %101 = OpCompositeExtract %_arr__arr_strided_arr_uint_2_uint_3 %value_2 0
- %98 = OpFunctionCall %void %assign_and_preserve_padding_2_s_a_X_el %100 %101
+ %99 = OpLabel
+ %101 = OpCompositeExtract %uint %dest_2 0
+ %102 = OpCompositeConstruct %_arr_uint_uint_1 %101
+ %103 = OpCompositeExtract %_arr__arr_strided_arr_uint_2_uint_3 %value_2 0
+ %100 = OpFunctionCall %void %assign_and_preserve_padding_2_s_a_X_el %102 %103
OpReturn
OpFunctionEnd
-%assign_and_preserve_padding_s_a = OpFunction %void None %102
+%assign_and_preserve_padding_s_a = OpFunction %void None %104
%value_3 = OpFunctionParameter %_arr_strided_arr_1_uint_4
- %105 = OpLabel
- %i_1 = OpVariable %_ptr_Function_uint Function %37
-%var_for_index_2 = OpVariable %_ptr_Function__arr_strided_arr_1_uint_4 Function %118
- OpStore %i_1 %37
- OpBranch %107
%107 = OpLabel
- OpLoopMerge %108 %109 None
- OpBranch %110
- %110 = OpLabel
- %112 = OpLoad %uint %i_1
- %113 = OpULessThan %bool %112 %uint_4
- %111 = OpLogicalNot %bool %113
- OpSelectionMerge %114 None
- OpBranchConditional %111 %115 %114
- %115 = OpLabel
- OpBranch %108
- %114 = OpLabel
- OpStore %var_for_index_2 %value_3
- %120 = OpLoad %uint %i_1
- %121 = OpCompositeConstruct %_arr_uint_uint_1 %120
- %122 = OpLoad %uint %i_1
- %124 = OpAccessChain %_ptr_Function_strided_arr_1 %var_for_index_2 %122
- %125 = OpLoad %strided_arr_1 %124
- %119 = OpFunctionCall %void %assign_and_preserve_padding_1_s_a_X %121 %125
+ %i_1 = OpVariable %_ptr_Function_uint Function %39
+%var_for_index_2 = OpVariable %_ptr_Function__arr_strided_arr_1_uint_4 Function %120
+ OpStore %i_1 %39
OpBranch %109
%109 = OpLabel
- %126 = OpLoad %uint %i_1
- %127 = OpIAdd %uint %126 %uint_1
- OpStore %i_1 %127
- OpBranch %107
- %108 = OpLabel
+ OpLoopMerge %110 %111 None
+ OpBranch %112
+ %112 = OpLabel
+ %114 = OpLoad %uint %i_1
+ %115 = OpULessThan %bool %114 %uint_4
+ %113 = OpLogicalNot %bool %115
+ OpSelectionMerge %116 None
+ OpBranchConditional %113 %117 %116
+ %117 = OpLabel
+ OpBranch %110
+ %116 = OpLabel
+ OpStore %var_for_index_2 %value_3
+ %122 = OpLoad %uint %i_1
+ %123 = OpCompositeConstruct %_arr_uint_uint_1 %122
+ %124 = OpLoad %uint %i_1
+ %126 = OpAccessChain %_ptr_Function_strided_arr_1 %var_for_index_2 %124
+ %127 = OpLoad %strided_arr_1 %126
+ %121 = OpFunctionCall %void %assign_and_preserve_padding_1_s_a_X %123 %127
+ OpBranch %111
+ %111 = OpLabel
+ %128 = OpLoad %uint %i_1
+ %129 = OpIAdd %uint %128 %uint_1
+ OpStore %i_1 %129
+ OpBranch %109
+ %110 = OpLabel
OpReturn
OpFunctionEnd
- %f_1 = OpFunction %void None %128
- %130 = OpLabel
- %132 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0 %uint_0
- %133 = OpLoad %_arr_strided_arr_1_uint_4 %132
- %136 = OpAccessChain %_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 %s %uint_0 %uint_0 %int_3 %uint_0
- %137 = OpLoad %_arr__arr_strided_arr_uint_2_uint_3 %136
- %140 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %s %uint_0 %uint_0 %int_3 %uint_0 %int_2
- %141 = OpLoad %_arr_strided_arr_uint_2 %140
+ %f_1 = OpFunction %void None %130
+ %132 = OpLabel
+ %134 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0 %uint_0
+ %135 = OpLoad %_arr_strided_arr_1_uint_4 %134
+ %138 = OpAccessChain %_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 %s %uint_0 %uint_0 %int_3 %uint_0
+ %139 = OpLoad %_arr__arr_strided_arr_uint_2_uint_3 %138
+ %141 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %s %uint_0 %uint_0 %int_3 %uint_0 %int_2
+ %142 = OpLoad %_arr_strided_arr_uint_2 %141
%143 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0
%144 = OpLoad %float %143
- %145 = OpFunctionCall %void %assign_and_preserve_padding_s_a %118
+ %145 = OpFunctionCall %void %assign_and_preserve_padding_s_a %120
%146 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0
OpStore %146 %float_5
OpReturn
OpFunctionEnd
- %f = OpFunction %void None %128
+ %f = OpFunction %void None %130
%149 = OpLabel
%150 = OpFunctionCall %void %f_1
OpReturn
diff --git a/test/tint/bug/tint/2059.wgsl.expected.glsl b/test/tint/bug/tint/2059.wgsl.expected.glsl
index 6a99d71..4b50ac1 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.glsl
+++ b/test/tint/bug/tint/2059.wgsl.expected.glsl
@@ -67,9 +67,9 @@
}
void assign_and_preserve_padding_buffer7_X_m_X(uint dest[2], mat3 value) {
- buffer7.inner[dest[0]].m[dest[0]][0] = value[0u];
- buffer7.inner[dest[0]].m[dest[0]][1] = value[1u];
- buffer7.inner[dest[0]].m[dest[0]][2] = value[2u];
+ buffer7.inner[dest[0]].m[dest[1]][0] = value[0u];
+ buffer7.inner[dest[0]].m[dest[1]][1] = value[1u];
+ buffer7.inner[dest[0]].m[dest[1]][2] = value[2u];
}
void assign_and_preserve_padding_buffer1_m(mat3 value) {
diff --git a/test/tint/bug/tint/2059.wgsl.expected.spvasm b/test/tint/bug/tint/2059.wgsl.expected.spvasm
index 4c40ec4..043a52c 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/2059.wgsl.expected.spvasm
@@ -286,17 +286,17 @@
%value_2 = OpFunctionParameter %mat3v3float
%89 = OpLabel
%90 = OpCompositeExtract %uint %dest_1 0
- %91 = OpCompositeExtract %uint %dest_1 0
+ %91 = OpCompositeExtract %uint %dest_1 1
%92 = OpAccessChain %_ptr_StorageBuffer_v3float %buffer7 %uint_0 %90 %uint_0 %91 %44
%93 = OpCompositeExtract %v3float %value_2 0
OpStore %92 %93
%94 = OpCompositeExtract %uint %dest_1 0
- %95 = OpCompositeExtract %uint %dest_1 0
+ %95 = OpCompositeExtract %uint %dest_1 1
%96 = OpAccessChain %_ptr_StorageBuffer_v3float %buffer7 %uint_0 %94 %uint_0 %95 %int_1
%97 = OpCompositeExtract %v3float %value_2 1
OpStore %96 %97
%98 = OpCompositeExtract %uint %dest_1 0
- %99 = OpCompositeExtract %uint %dest_1 0
+ %99 = OpCompositeExtract %uint %dest_1 1
%100 = OpAccessChain %_ptr_StorageBuffer_v3float %buffer7 %uint_0 %98 %uint_0 %99 %int_2
%101 = OpCompositeExtract %v3float %value_2 2
OpStore %100 %101
