[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);
}
)");
}
