Google Git
Sign in
dawn/dawn/refs/heads/chromium/6873^!/.
commit
12a3b24c456cebd9fd11f23ac0164f78129b00c6
[log]
author
Ryan Harrison <rharrison@chromium.org>
Wed Nov 27 22:01:26 2024 +0000
committer
Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com>
Wed Nov 27 22:01:26 2024 +0000
tree
299b18e83c708bf078d274dd89bbbe51397b17f9
parent
031b57680c4d7c48de49c3d1f44c054879ddc17d [diff]
[tint][ir][val] Enforce rules around scope for function vars

- Adds a check that function address space vars are only declared in a
  function scope
- Adds a check that non-function address space are not declared in a
  function scope
- Add a capability to allow MSL specific IR to have private address
  space vars declared in a function scope.
- Bunch of test fixes/updates

Fixes: 381109471
Fixes: 381109773
Change-Id: I874930a858777c43a0c26354c7d888ffeab9056c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/216957
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

diff --git a/src/tint/lang/core/ir/transform/remove_terminator_args.h b/src/tint/lang/core/ir/transform/remove_terminator_args.h
index 0b366b8..63bb87d 100644
--- a/src/tint/lang/core/ir/transform/remove_terminator_args.h
+++ b/src/tint/lang/core/ir/transform/remove_terminator_args.h

@@ -45,6 +45,7 @@
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,
     core::ir::Capability::kAllowClipDistancesOnF32,
+    core::ir::Capability::kAllowPrivateVarsInFunctions,
 };
 
 /// RemoveTerminatorArgs is a transform that removes all arguments from terminator instructions and

diff --git a/src/tint/lang/core/ir/transform/rename_conflicts.h b/src/tint/lang/core/ir/transform/rename_conflicts.h
index a56cab3..fb5d4b5 100644
--- a/src/tint/lang/core/ir/transform/rename_conflicts.h
+++ b/src/tint/lang/core/ir/transform/rename_conflicts.h

@@ -45,6 +45,7 @@
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,
     core::ir::Capability::kAllowClipDistancesOnF32,
+    core::ir::Capability::kAllowPrivateVarsInFunctions,
 };
 
 /// RenameConflicts is a transform that renames declarations which prevent identifiers from

diff --git a/src/tint/lang/core/ir/transform/rename_conflicts_test.cc b/src/tint/lang/core/ir/transform/rename_conflicts_test.cc
index e418b0c..b1be9ad 100644
--- a/src/tint/lang/core/ir/transform/rename_conflicts_test.cc
+++ b/src/tint/lang/core/ir/transform/rename_conflicts_test.cc

@@ -124,7 +124,7 @@
 
 TEST_F(IRToProgramRenameConflictsTest, Conflict_RootBlockVarAndStructWithSameName) {
     auto* s = ty.Struct(b.ir.symbols.New("v"), {{b.ir.symbols.New("x"), ty.i32()}});
-    b.Append(mod.root_block, [&] { b.ir.SetName(b.Var(ty.ptr(function, s)), "v"); });
+    b.Append(mod.root_block, [&] { b.ir.SetName(b.Var(ty.ptr(private_, s)), "v"); });
 
     auto* src = R"(
 v = struct @align(4) {
@@ -132,7 +132,7 @@
 }
 
 $B1: {  # root
-  %v:ptr<function, v, read_write> = var
+  %v:ptr<private, v, read_write> = var
 }
 
 )";
@@ -144,7 +144,7 @@
 }
 
 $B1: {  # root
-  %v_1:ptr<function, v, read_write> = var
+  %v_1:ptr<private, v, read_write> = var
 }
 
 )";

diff --git a/src/tint/lang/core/ir/transform/value_to_let.h b/src/tint/lang/core/ir/transform/value_to_let.h
index 945f773..852c554 100644
--- a/src/tint/lang/core/ir/transform/value_to_let.h
+++ b/src/tint/lang/core/ir/transform/value_to_let.h

@@ -46,6 +46,7 @@
     core::ir::Capability::kAllowVectorElementPointer,
     core::ir::Capability::kAllowHandleVarsWithoutBindings,
     core::ir::Capability::kAllowClipDistancesOnF32,
+    core::ir::Capability::kAllowPrivateVarsInFunctions,
 };
 
 /// Configuration for ValueToLet transform.

diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 7677f4c..70fc255 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc

@@ -2377,14 +2377,25 @@
         return;
     }
 
-    {
-        auto result = ValidateBindingPoint(var->BindingPoint(), mv->AddressSpace());
-        if (result != Success) {
-            AddError(var) << result.Failure();
+    if (var->Block() != mod_.root_block && mv->AddressSpace() != AddressSpace::kFunction) {
+        if (!capabilities_.Contains(Capability::kAllowPrivateVarsInFunctions) ||
+            mv->AddressSpace() != AddressSpace::kPrivate) {
+            AddError(var) << "vars in a function scope must be in the 'function' address space";
             return;
         }
     }
 
+    if (var->Block() == mod_.root_block && mv->AddressSpace() == AddressSpace::kFunction) {
+        AddError(var) << "vars in the 'function' address space must be in a function scope";
+        return;
+    }
+
+    if (auto result = ValidateBindingPoint(var->BindingPoint(), mv->AddressSpace());
+        result != Success) {
+        AddError(var) << result.Failure();
+        return;
+    }
+
     if (mv->AddressSpace() == AddressSpace::kWorkgroup) {
         if (auto* ary = result_type->UnwrapPtr()->As<core::type::Array>()) {
             if (auto* count = ary->Count()->As<core::ir::type::ValueArrayCount>()) {

diff --git a/src/tint/lang/core/ir/validator.h b/src/tint/lang/core/ir/validator.h
index cf4809a..dcf8564 100644
--- a/src/tint/lang/core/ir/validator.h
+++ b/src/tint/lang/core/ir/validator.h

@@ -58,6 +58,8 @@
     kAllowRefTypes,
     /// Allows access instructions to create pointers to vector elements.
     kAllowVectorElementPointer,
+    /// Allows private address space variables in function scopes.
+    kAllowPrivateVarsInFunctions,
 };
 
 /// Capabilities is a set of Capability

diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 57f72f0..a9dfb11 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc

@@ -1552,7 +1552,7 @@
 
     b.Append(f->Block(), [&] {
         auto* l = b.Load(invalid);
-        auto* v = b.Var("v", AddressSpace::kPrivate, ty.bool_());
+        auto* v = b.Var("v", AddressSpace::kFunction, ty.bool_());
         v->SetInitializer(l->Result(0));
         b.Unreachable();
     });
@@ -1573,7 +1573,7 @@
 %f = @fragment func():void {
   $B2: {
     %3:bool = load %invalid
-    %v:ptr<private, bool, read_write> = var, %3
+    %v:ptr<function, bool, read_write> = var, %3
     unreachable
   }
 }
@@ -5150,21 +5150,42 @@
 )");
 }
 
-TEST_F(IR_ValidatorTest, Var_Private_UnexpectedInputAttachmentIndex) {
+TEST_F(IR_ValidatorTest, Var_Function_OutsideFunctionScope) {
+    auto* v = b.Var<function, f32>();
+    mod.root_block->Append(v);
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:2:39 error: var: vars in the 'function' address space must be in a function scope
+  %1:ptr<function, f32, read_write> = var
+                                      ^^^
+
+:1:1 note: in block
+$B1: {  # root
+^^^
+
+note: # Disassembly
+$B1: {  # root
+  %1:ptr<function, f32, read_write> = var
+}
+
+)");
+}
+
+TEST_F(IR_ValidatorTest, Var_NonFunction_InsideFunctionScope) {
     auto* f = b.Function("my_func", ty.void_());
 
     b.Append(f->Block(), [&] {
-        auto* v = b.Var<private_, f32>();
-
-        v->SetInputAttachmentIndex(0);
+        b.Var<private_, f32>();
         b.Return(f);
     });
 
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(:3:40 error: var: '@input_attachment_index' is not valid for non-handle var
-    %2:ptr<private, f32, read_write> = var @input_attachment_index(0)
+              R"(:3:40 error: var: vars in a function scope must be in the 'function' address space
+    %2:ptr<private, f32, read_write> = var
                                        ^^^
 
 :2:3 note: in block
@@ -5174,132 +5195,144 @@
 note: # Disassembly
 %my_func = func():void {
   $B1: {
-    %2:ptr<private, f32, read_write> = var @input_attachment_index(0)
+    %2:ptr<private, f32, read_write> = var
     ret
   }
 }
 )");
 }
 
+TEST_F(IR_ValidatorTest, Var_Private_InsideFunctionScopeWithCapability) {
+    auto* f = b.Function("my_func", ty.void_());
+
+    b.Append(f->Block(), [&] {
+        b.Var<private_, f32>();
+        b.Return(f);
+    });
+
+    auto res = ir::Validate(mod, Capabilities{Capability::kAllowPrivateVarsInFunctions});
+    ASSERT_EQ(res, Success);
+}
+
+TEST_F(IR_ValidatorTest, Var_Private_UnexpectedInputAttachmentIndex) {
+    auto* v = b.Var<private_, f32>();
+    v->SetInputAttachmentIndex(0);
+    mod.root_block->Append(v);
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:2:38 error: var: '@input_attachment_index' is not valid for non-handle var
+  %1:ptr<private, f32, read_write> = var @input_attachment_index(0)
+                                     ^^^
+
+:1:1 note: in block
+$B1: {  # root
+^^^
+
+note: # Disassembly
+$B1: {  # root
+  %1:ptr<private, f32, read_write> = var @input_attachment_index(0)
+}
+
+)");
+}
+
 TEST_F(IR_ValidatorTest, Var_PushConstant_UnexpectedInputAttachmentIndex) {
-    auto* f = b.Function("my_func", ty.void_());
-
-    b.Append(f->Block(), [&] {
-        auto* v = b.Var<push_constant, f32>();
-        v->SetInputAttachmentIndex(0);
-        b.Return(f);
-    });
+    auto* v = b.Var<push_constant, f32>();
+    v->SetInputAttachmentIndex(0);
+    mod.root_block->Append(v);
 
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(:3:40 error: var: '@input_attachment_index' is not valid for non-handle var
-    %2:ptr<push_constant, f32, read> = var @input_attachment_index(0)
-                                       ^^^
+              R"(:2:38 error: var: '@input_attachment_index' is not valid for non-handle var
+  %1:ptr<push_constant, f32, read> = var @input_attachment_index(0)
+                                     ^^^
 
-:2:3 note: in block
-  $B1: {
-  ^^^
+:1:1 note: in block
+$B1: {  # root
+^^^
 
 note: # Disassembly
-%my_func = func():void {
-  $B1: {
-    %2:ptr<push_constant, f32, read> = var @input_attachment_index(0)
-    ret
-  }
+$B1: {  # root
+  %1:ptr<push_constant, f32, read> = var @input_attachment_index(0)
 }
+
 )");
 }
 
 TEST_F(IR_ValidatorTest, Var_Storage_UnexpectedInputAttachmentIndex) {
-    auto* f = b.Function("my_func", ty.void_());
-
-    b.Append(f->Block(), [&] {
-        auto* v = b.Var<storage, f32>();
-        v->SetBindingPoint(0, 0);
-        v->SetInputAttachmentIndex(0);
-        b.Return(f);
-    });
+    auto* v = b.Var<storage, f32>();
+    v->SetBindingPoint(0, 0);
+    v->SetInputAttachmentIndex(0);
+    mod.root_block->Append(v);
 
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(:3:40 error: var: '@input_attachment_index' is not valid for non-handle var
-    %2:ptr<storage, f32, read_write> = var @binding_point(0, 0) @input_attachment_index(0)
-                                       ^^^
+              R"(:2:38 error: var: '@input_attachment_index' is not valid for non-handle var
+  %1:ptr<storage, f32, read_write> = var @binding_point(0, 0) @input_attachment_index(0)
+                                     ^^^
 
-:2:3 note: in block
-  $B1: {
-  ^^^
+:1:1 note: in block
+$B1: {  # root
+^^^
 
 note: # Disassembly
-%my_func = func():void {
-  $B1: {
-    %2:ptr<storage, f32, read_write> = var @binding_point(0, 0) @input_attachment_index(0)
-    ret
-  }
+$B1: {  # root
+  %1:ptr<storage, f32, read_write> = var @binding_point(0, 0) @input_attachment_index(0)
 }
+
 )");
 }
 
 TEST_F(IR_ValidatorTest, Var_Uniform_UnexpectedInputAttachmentIndex) {
-    auto* f = b.Function("my_func", ty.void_());
-
-    b.Append(f->Block(), [&] {
-        auto* v = b.Var<uniform, f32>();
-        v->SetBindingPoint(0, 0);
-        v->SetInputAttachmentIndex(0);
-        b.Return(f);
-    });
+    auto* v = b.Var<uniform, f32>();
+    v->SetBindingPoint(0, 0);
+    v->SetInputAttachmentIndex(0);
+    mod.root_block->Append(v);
 
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(:3:34 error: var: '@input_attachment_index' is not valid for non-handle var
-    %2:ptr<uniform, f32, read> = var @binding_point(0, 0) @input_attachment_index(0)
-                                 ^^^
+              R"(:2:32 error: var: '@input_attachment_index' is not valid for non-handle var
+  %1:ptr<uniform, f32, read> = var @binding_point(0, 0) @input_attachment_index(0)
+                               ^^^
 
-:2:3 note: in block
-  $B1: {
-  ^^^
+:1:1 note: in block
+$B1: {  # root
+^^^
 
 note: # Disassembly
-%my_func = func():void {
-  $B1: {
-    %2:ptr<uniform, f32, read> = var @binding_point(0, 0) @input_attachment_index(0)
-    ret
-  }
+$B1: {  # root
+  %1:ptr<uniform, f32, read> = var @binding_point(0, 0) @input_attachment_index(0)
 }
+
 )");
 }
 
 TEST_F(IR_ValidatorTest, Var_Workgroup_UnexpectedInputAttachmentIndex) {
-    auto* f = b.Function("my_func", ty.void_());
-
-    b.Append(f->Block(), [&] {
-        auto* v = b.Var<workgroup, f32>();
-        v->SetInputAttachmentIndex(0);
-        b.Return(f);
-    });
+    auto* v = b.Var<workgroup, f32>();
+    v->SetInputAttachmentIndex(0);
+    mod.root_block->Append(v);
 
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(:3:42 error: var: '@input_attachment_index' is not valid for non-handle var
-    %2:ptr<workgroup, f32, read_write> = var @input_attachment_index(0)
-                                         ^^^
+              R"(:2:40 error: var: '@input_attachment_index' is not valid for non-handle var
+  %1:ptr<workgroup, f32, read_write> = var @input_attachment_index(0)
+                                       ^^^
 
-:2:3 note: in block
-  $B1: {
-  ^^^
+:1:1 note: in block
+$B1: {  # root
+^^^
 
 note: # Disassembly
-%my_func = func():void {
-  $B1: {
-    %2:ptr<workgroup, f32, read_write> = var @input_attachment_index(0)
-    ret
-  }
+$B1: {  # root
+  %1:ptr<workgroup, f32, read_write> = var @input_attachment_index(0)
 }
+
 )");
 }
 

diff --git a/src/tint/lang/glsl/writer/type_test.cc b/src/tint/lang/glsl/writer/type_test.cc
index 543d05d..487394f 100644
--- a/src/tint/lang/glsl/writer/type_test.cc
+++ b/src/tint/lang/glsl/writer/type_test.cc

@@ -47,7 +47,7 @@
 TEST_F(GlslWriterTest, EmitType_Array) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array<bool, 4>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.array<bool, 4>()));
         b.Return(func);
     });
 
@@ -63,7 +63,7 @@
 TEST_F(GlslWriterTest, EmitType_ArrayOfArray) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array(ty.array<bool, 4>(), 5)));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.array(ty.array<bool, 4>(), 5)));
         b.Return(func);
     });
 
@@ -80,7 +80,7 @@
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
         b.Var("a",
-              ty.ptr(core::AddressSpace::kPrivate, ty.array(ty.array(ty.array<bool, 4>(), 5), 6)));
+              ty.ptr(core::AddressSpace::kFunction, ty.array(ty.array(ty.array<bool, 4>(), 5), 6)));
         b.Return(func);
     });
 
@@ -100,7 +100,7 @@
                                             });
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, Inner));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, Inner));
         b.Return(func);
     });
 
@@ -121,7 +121,7 @@
 TEST_F(GlslWriterTest, EmitType_Bool) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.bool_()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.bool_()));
         b.Return(func);
     });
 
@@ -137,7 +137,7 @@
 TEST_F(GlslWriterTest, EmitType_F32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.f32()));
         b.Return(func);
     });
 
@@ -153,7 +153,7 @@
 TEST_F(GlslWriterTest, EmitType_F16) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f16()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.f16()));
         b.Return(func);
     });
 
@@ -170,7 +170,7 @@
 TEST_F(GlslWriterTest, EmitType_I32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.i32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.i32()));
         b.Return(func);
     });
 
@@ -186,7 +186,7 @@
 TEST_F(GlslWriterTest, EmitType_Matrix_F32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.mat2x3<f32>()));
         b.Return(func);
     });
 
@@ -202,7 +202,7 @@
 TEST_F(GlslWriterTest, EmitType_MatrixSquare_F32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x2<f32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.mat2x2<f32>()));
         b.Return(func);
     });
 
@@ -218,7 +218,7 @@
 TEST_F(GlslWriterTest, EmitType_Matrix_F16) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f16>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.mat2x3<f16>()));
         b.Return(func);
     });
 
@@ -235,7 +235,7 @@
 TEST_F(GlslWriterTest, EmitType_U32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.u32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.u32()));
         b.Return(func);
     });
 
@@ -277,7 +277,7 @@
 TEST_F(GlslWriterTest, EmitType_Vector_F32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec3<f32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec3<f32>()));
         b.Return(func);
     });
 
@@ -293,7 +293,7 @@
 TEST_F(GlslWriterTest, EmitType_Vector_F16) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec3<f16>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec3<f16>()));
         b.Return(func);
     });
 
@@ -310,7 +310,7 @@
 TEST_F(GlslWriterTest, EmitType_Vector_I32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec2<i32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec2<i32>()));
         b.Return(func);
     });
 
@@ -326,7 +326,7 @@
 TEST_F(GlslWriterTest, EmitType_Vector_U32) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec4<u32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec4<u32>()));
         b.Return(func);
     });
 
@@ -342,7 +342,7 @@
 TEST_F(GlslWriterTest, EmitType_Vector_bool) {
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec3<bool>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec3<bool>()));
         b.Return(func);
     });
 
@@ -375,7 +375,7 @@
                                               });
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, s));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, s));
         b.Return(func);
     });
 
@@ -401,8 +401,8 @@
                                               });
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, s));
-        b.Var("b", ty.ptr(core::AddressSpace::kPrivate, s));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, s));
+        b.Var("b", ty.ptr(core::AddressSpace::kFunction, s));
         b.Return(func);
     });
 
@@ -435,7 +435,7 @@
                                               });
     auto* func = b.ComputeFunction("foo");
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, s));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, s));
         b.Return(func);
     });
 

diff --git a/src/tint/lang/hlsl/writer/raise/promote_initializers_test.cc b/src/tint/lang/hlsl/writer/raise/promote_initializers_test.cc
index a9f978d..e63a111 100644
--- a/src/tint/lang/hlsl/writer/raise/promote_initializers_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/promote_initializers_test.cc

@@ -43,14 +43,14 @@
 TEST_F(HlslWriterPromoteInitializersTest, NoStructInitializers) {
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     b.Append(func->Block(), [&] {
-        b.Var<private_>("a", b.Zero<i32>());
+        b.Var<function>("a", b.Zero<i32>());
         b.Return(func);
     });
 
     auto* src = R"(
 %foo = @fragment func():void {
   $B1: {
-    %a:ptr<private, i32, read_write> = var, 0i
+    %a:ptr<function, i32, read_write> = var, 0i
     ret
   }
 }
@@ -70,7 +70,7 @@
 
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     b.Append(func->Block(), [&] {
-        b.Var<private_>("a", b.Composite(str_ty, 1_i));
+        b.Var<function>("a", b.Composite(str_ty, 1_i));
         b.Return(func);
     });
 
@@ -81,7 +81,7 @@
 
 %foo = @fragment func():void {
   $B1: {
-    %a:ptr<private, S, read_write> = var, S(1i)
+    %a:ptr<function, S, read_write> = var, S(1i)
     ret
   }
 }
@@ -97,14 +97,14 @@
 TEST_F(HlslWriterPromoteInitializersTest, ArrayInVarNoChange) {
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     b.Append(func->Block(), [&] {
-        b.Var<private_>("a", b.Zero<array<i32, 2>>());
+        b.Var<function>("a", b.Zero<array<i32, 2>>());
         b.Return(func);
     });
 
     auto* src = R"(
 %foo = @fragment func():void {
   $B1: {
-    %a:ptr<private, array<i32, 2>, read_write> = var, array<i32, 2>(0i)
+    %a:ptr<function, array<i32, 2>, read_write> = var, array<i32, 2>(0i)
     ret
   }
 }

diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 5bb79e2..a2a1e21 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc

@@ -121,12 +121,13 @@
 
     /// @returns the generated MSL shader
     tint::Result<PrintResult> Generate() {
-        auto valid =
-            core::ir::ValidateAndDumpIfNeeded(ir_, "msl.Printer",
-                                              core::ir::Capabilities{
-                                                  core::ir::Capability::kAllow8BitIntegers,
-                                                  core::ir::Capability::kAllowPointersInStructures,
-                                              });
+        auto valid = core::ir::ValidateAndDumpIfNeeded(
+            ir_, "msl.Printer",
+            core::ir::Capabilities{
+                core::ir::Capability::kAllow8BitIntegers,
+                core::ir::Capability::kAllowPointersInStructures,
+                core::ir::Capability::kAllowPrivateVarsInFunctions,
+            });
         if (valid != Success) {
             return std::move(valid.Failure());
         }

diff --git a/src/tint/lang/msl/writer/raise/binary_polyfill.cc b/src/tint/lang/msl/writer/raise/binary_polyfill.cc
index bdf4b07..536065a 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill.cc

@@ -160,6 +160,7 @@
     auto result = ValidateAndDumpIfNeeded(ir, "msl.BinaryPolyfill",
                                           core::ir::Capabilities{
                                               core::ir::Capability::kAllowPointersInStructures,
+                                              core::ir::Capability::kAllowPrivateVarsInFunctions,
                                           });
     if (result != Success) {
         return result.Failure();

diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 8f3983a..cca1607 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc

@@ -936,6 +936,7 @@
     auto result = ValidateAndDumpIfNeeded(ir, "msl.BuiltinPolyfill",
                                           core::ir::Capabilities{
                                               core::ir::Capability::kAllowPointersInStructures,
+                                              core::ir::Capability::kAllowPrivateVarsInFunctions,
                                           });
     if (result != Success) {
         return result.Failure();

diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.h b/src/tint/lang/msl/writer/raise/module_scope_vars.h
index a3c69ee..f9570e7 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.h
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.h

@@ -30,6 +30,7 @@
 
 #include <string>
 
+#include "src/tint/lang/core/ir/validator.h"
 #include "src/tint/utils/result/result.h"
 
 // Forward declarations.

diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 775b6f7..7a64b1d 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc

@@ -41,7 +41,10 @@
 
 class MslWriter_ModuleScopeVarsTest : public core::ir::transform::TransformTest {
   public:
-    void SetUp() override { capabilities.Add(core::ir::Capability::kAllowPointersInStructures); }
+    void SetUp() override {
+        capabilities.Add(core::ir::Capability::kAllowPointersInStructures,
+                         core::ir::Capability::kAllowPrivateVarsInFunctions);
+    }
 };
 
 TEST_F(MslWriter_ModuleScopeVarsTest, NoModuleScopeVars) {

diff --git a/src/tint/lang/msl/writer/raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/raise/packed_vec3_test.cc
index a230647..d65648f 100644
--- a/src/tint/lang/msl/writer/raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/raise/packed_vec3_test.cc

@@ -41,7 +41,7 @@
 using MslWriter_PackedVec3Test = core::ir::transform::TransformTest;
 
 TEST_F(MslWriter_PackedVec3Test, NoModify_PrivateVar) {
-    auto* var = b.Var<function, vec3<u32>>("v");
+    auto* var = b.Var<private_, vec3<u32>>("v");
     mod.root_block->Append(var);
 
     auto* func = b.Function("foo", ty.vec3<u32>());
@@ -51,7 +51,7 @@
 
     auto* src = R"(
 $B1: {  # root
-  %v:ptr<function, vec3<u32>, read_write> = var
+  %v:ptr<private, vec3<u32>, read_write> = var
 }
 
 %foo = func():vec3<u32> {

diff --git a/src/tint/lang/msl/writer/raise/unary_polyfill.cc b/src/tint/lang/msl/writer/raise/unary_polyfill.cc
index 940f16a..072bb95 100644
--- a/src/tint/lang/msl/writer/raise/unary_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/unary_polyfill.cc

@@ -91,6 +91,7 @@
     auto result = ValidateAndDumpIfNeeded(ir, "msl.UnaryPolyfill",
                                           core::ir::Capabilities{
                                               core::ir::Capability::kAllowPointersInStructures,
+                                              core::ir::Capability::kAllowPrivateVarsInFunctions,
                                           });
     if (result != Success) {
         return result.Failure();

diff --git a/src/tint/lang/msl/writer/type_test.cc b/src/tint/lang/msl/writer/type_test.cc
index 502ae94..9c41fdf 100644
--- a/src/tint/lang/msl/writer/type_test.cc
+++ b/src/tint/lang/msl/writer/type_test.cc

@@ -47,14 +47,14 @@
 TEST_F(MslWriterTest, EmitType_Array) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array<bool, 4>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.array<bool, 4>()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
-  thread tint_array<bool, 4> a = {};
+  tint_array<bool, 4> a = {};
 }
 )");
 }
@@ -62,14 +62,14 @@
 TEST_F(MslWriterTest, EmitType_ArrayOfArray) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array(ty.array<bool, 4>(), 5)));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.array(ty.array<bool, 4>(), 5)));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
-  thread tint_array<tint_array<bool, 4>, 5> a = {};
+  tint_array<tint_array<bool, 4>, 5> a = {};
 }
 )");
 }
@@ -78,14 +78,14 @@
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a",
-              ty.ptr(core::AddressSpace::kPrivate, ty.array(ty.array(ty.array<bool, 4>(), 5), 6)));
+              ty.ptr(core::AddressSpace::kFunction, ty.array(ty.array(ty.array<bool, 4>(), 5), 6)));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
-  thread tint_array<tint_array<tint_array<bool, 4>, 5>, 6> a = {};
+  tint_array<tint_array<tint_array<bool, 4>, 5>, 6> a = {};
 }
 )");
 }
@@ -93,14 +93,14 @@
 TEST_F(MslWriterTest, EmitType_RuntimeArray) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array<bool, 0>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.array<bool, 0>()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
-  thread tint_array<bool, 1> a = {};
+  tint_array<bool, 1> a = {};
 }
 )");
 }
@@ -108,14 +108,14 @@
 TEST_F(MslWriterTest, EmitType_Bool) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.bool_()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.bool_()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread bool a = false;
+  bool a = false;
 }
 )");
 }
@@ -123,14 +123,14 @@
 TEST_F(MslWriterTest, EmitType_F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.f32()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread float a = 0.0f;
+  float a = 0.0f;
 }
 )");
 }
@@ -138,14 +138,14 @@
 TEST_F(MslWriterTest, EmitType_F16) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f16()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.f16()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread half a = 0.0h;
+  half a = 0.0h;
 }
 )");
 }
@@ -153,14 +153,14 @@
 TEST_F(MslWriterTest, EmitType_I32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.i32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.i32()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread int a = 0;
+  int a = 0;
 }
 )");
 }
@@ -168,14 +168,14 @@
 TEST_F(MslWriterTest, EmitType_Matrix_F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.mat2x3<f32>()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread float2x3 a = float2x3(0.0f);
+  float2x3 a = float2x3(0.0f);
 }
 )");
 }
@@ -183,28 +183,28 @@
 TEST_F(MslWriterTest, EmitType_Matrix_F16) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f16>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.mat2x3<f16>()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread half2x3 a = half2x3(0.0h);
+  half2x3 a = half2x3(0.0h);
 }
 )");
 }
 TEST_F(MslWriterTest, EmitType_U32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.u32()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.u32()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread uint a = 0u;
+  uint a = 0u;
 }
 )");
 }
@@ -242,14 +242,14 @@
 TEST_F(MslWriterTest, EmitType_Vector) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec3<f32>()));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.vec3<f32>()));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread float3 a = 0.0f;
+  float3 a = 0.0f;
 }
 )");
 }
@@ -320,7 +320,7 @@
                                               });
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, s));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, s));
         b.Return(func);
     });
 
@@ -332,7 +332,7 @@
 };
 
 void foo() {
-  thread S a = {};
+  S a = {};
 }
 )");
 }
@@ -344,8 +344,8 @@
                                               });
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, s));
-        b.Var("b", ty.ptr(core::AddressSpace::kPrivate, s));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, s));
+        b.Var("b", ty.ptr(core::AddressSpace::kFunction, s));
         b.Return(func);
     });
 
@@ -357,8 +357,8 @@
 };
 
 void foo() {
-  thread S a = {};
-  thread S b = {};
+  S a = {};
+  S b = {};
 }
 )");
 }
@@ -1145,14 +1145,14 @@
 TEST_F(MslWriterTest, EmitType_SubgroupMatrixLeft) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.subgroup_matrix_left(ty.f32(), 8, 8)));
+        b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.subgroup_matrix_left(ty.f32(), 8, 8)));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate({}, validate::MslVersion::kMsl_2_3)) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread simdgroup_float8x8 a = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+  simdgroup_float8x8 a = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
 }
 )");
 }
@@ -1177,14 +1177,15 @@
 TEST_F(MslWriterTest, EmitType_SubgroupMatrixResult) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
-        b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.subgroup_matrix_result(ty.f32(), 8, 8)));
+        b.Var("a",
+              ty.ptr(core::AddressSpace::kFunction, ty.subgroup_matrix_result(ty.f32(), 8, 8)));
         b.Return(func);
     });
 
     ASSERT_TRUE(Generate({}, validate::MslVersion::kMsl_2_3)) << err_ << output_.msl;
     EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
-  thread simdgroup_float8x8 a = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+  simdgroup_float8x8 a = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
 }
 )");
 }