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) {