[tint][ir][val] Improve checks regarding binding points
- Refactors related code into common utility function
- Requires binding points for resource variables/param
- Rejects binding points on non-resource variables/params
- Reject binding point on non-entry point input parameters
Fixes: 377850460
Change-Id: I041088d42227f7344101bfffd9a52477509907ad
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/214354
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
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 3cca4ac..d5ded41 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
@@ -2531,7 +2531,6 @@
b.Append(b.ir.root_block,
[&] { //
input = b.Var("U", ty.ptr<workgroup>(T));
- input->SetBindingPoint(0, 0);
});
auto* f2 = b.Function("f2", T3);
@@ -2566,7 +2565,7 @@
auto* src = R"(
$B1: { # root
- %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+ %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
}
%f2 = func(%p:ptr<workgroup, array<vec4<i32>, 5>, read_write>):vec4<i32> {
@@ -2602,7 +2601,7 @@
auto* expect = R"(
$B1: { # root
- %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var @binding_point(0, 0)
+ %U:ptr<workgroup, array<array<array<vec4<i32>, 5>, 5>, 5>, read_write> = var
}
%f2 = func(%p_indices:array<u32, 2>):vec4<i32> {
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 3abb1dc..7655fbc 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -992,6 +992,17 @@
/// @param var the var to validate
void CheckVar(const Var* var);
+ /// Validates binding_point usage for pointers
+ /// @param binding_point the binding information associated with pointer
+ /// @param address_space the address space of pointer
+ /// @param target_str string to insert in error message describing what has a binding_point,
+ /// defaults to 'variable'
+ /// @returns Success if a valid usage, or reason for invalidity in Failure
+ Result<SuccessType, std::string> ValidateBindingPoint(
+ const std::optional<struct BindingPoint>& binding_point,
+ AddressSpace address_space,
+ const std::string& target_str = "variable");
+
/// Validates the given let
/// @param l the let to validate
void CheckLet(const Let* l);
@@ -1884,6 +1895,29 @@
"entry point params can only be a bool for fragment shaders"));
}
+ AddressSpace address_space = AddressSpace::kUndefined;
+ auto* mv = param->Type()->As<type::MemoryView>();
+ if (mv) {
+ address_space = mv->AddressSpace();
+ } else {
+ // ModuleScopeVars transform in MSL backends unwraps pointers to handles
+ if (param->Type()->IsAnyOf<type::Texture, type::Sampler>()) {
+ address_space = AddressSpace::kHandle;
+ }
+ }
+
+ if (func->Stage() != Function::PipelineStage::kUndefined) {
+ auto result = ValidateBindingPoint(param->BindingPoint(), address_space, "input param");
+ if (result != Success) {
+ AddError(param) << result.Failure();
+ }
+ } else {
+ if (param->BindingPoint().has_value()) {
+ AddError(param)
+ << "input param to non-entry point function has a binding point set";
+ }
+ }
+
scope_stack_.Add(param);
}
@@ -2275,23 +2309,12 @@
return;
}
- // Check that only resource variables have @group and @binding set
- switch (mv->AddressSpace()) {
- case AddressSpace::kHandle:
- if (!capabilities_.Contains(Capability::kAllowHandleVarsWithoutBindings)) {
- if (!var->BindingPoint().has_value()) {
- AddError(var) << "resource variable missing binding points";
- }
- }
- break;
- case AddressSpace::kStorage:
- case AddressSpace::kUniform:
- if (!var->BindingPoint().has_value()) {
- AddError(var) << "resource variable missing binding points";
- }
- break;
- default:
- break;
+ {
+ auto result = ValidateBindingPoint(var->BindingPoint(), mv->AddressSpace());
+ if (result != Success) {
+ AddError(var) << result.Failure();
+ return;
+ }
}
// Check that non-handle variables don't have @input_attachment_index set
@@ -2321,6 +2344,33 @@
}
}
+Result<SuccessType, std::string> Validator::ValidateBindingPoint(
+ const std::optional<struct BindingPoint>& binding_point,
+ AddressSpace address_space,
+ const std::string& target_str) {
+ switch (address_space) {
+ case AddressSpace::kHandle:
+ if (!capabilities_.Contains(Capability::kAllowHandleVarsWithoutBindings)) {
+ if (!binding_point.has_value()) {
+ return "a resource " + target_str + " is missing binding point";
+ }
+ }
+ break;
+ case AddressSpace::kStorage:
+ case AddressSpace::kUniform:
+ if (!binding_point.has_value()) {
+ return "a resource " + target_str + " is missing binding point";
+ }
+ break;
+ default:
+ if (binding_point.has_value()) {
+ return "a non-resource " + target_str + " has binding point";
+ }
+ break;
+ }
+ return Success;
+}
+
void Validator::CheckLet(const Let* l) {
if (!CheckResultsAndOperands(l, Let::kNumResults, Let::kNumOperands)) {
return;
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 81bc844..55cb1d2 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -661,6 +661,30 @@
)");
}
+TEST_F(IR_ValidatorTest, Function_Param_BindingPointWithoutCapability) {
+ auto* f = b.Function("my_func", ty.void_());
+ auto* p = b.FunctionParam("my_param", ty.ptr<uniform, i32>());
+ p->SetBindingPoint(0, 0);
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(res.Failure().reason.Str(),
+ R"(:1:17 error: input param to non-entry point function has a binding point set
+%my_func = func(%my_param:ptr<uniform, i32, read> [@binding_point(0, 0)]):void {
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:ptr<uniform, i32, read> [@binding_point(0, 0)]):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, Function_Return_BothLocationAndBuiltin) {
auto* f = VertexEntryPoint("my_func");
IOAttributes attr;
@@ -5089,7 +5113,7 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:2:31 error: var: resource variable missing binding points
+ R"(:2:31 error: var: a resource variable is missing binding point
%1:ptr<handle, i32, read> = var
^^^
@@ -5112,7 +5136,7 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:2:38 error: var: resource variable missing binding points
+ R"(:2:38 error: var: a resource variable is missing binding point
%1:ptr<storage, i32, read_write> = var
^^^
@@ -5135,7 +5159,7 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:2:32 error: var: resource variable missing binding points
+ R"(:2:32 error: var: a resource variable is missing binding point
%1:ptr<uniform, i32, read> = var
^^^
@@ -5151,13 +5175,36 @@
)");
}
+TEST_F(IR_ValidatorTest, Var_NonResourceWithBindingPoint) {
+ auto* v = b.Var(ty.ptr<private_, i32>());
+ v->SetBindingPoint(0, 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: a non-resource variable has binding point
+ %1:ptr<private, i32, read_write> = var @binding_point(0, 0)
+ ^^^
+
+:1:1 note: in block
+$B1: { # root
+^^^
+
+note: # Disassembly
+$B1: { # root
+ %1:ptr<private, i32, read_write> = var @binding_point(0, 0)
+}
+
+)");
+}
+
TEST_F(IR_ValidatorTest, Var_IOBothLocationAndBuiltin) {
auto* v = b.Var<AddressSpace::kIn, vec4<f32>>();
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
v->SetAttributes(attr);
- v->SetBindingPoint(0, 0);
mod.root_block->Append(v);
auto res = ir::Validate(mod);
@@ -5165,7 +5212,7 @@
EXPECT_EQ(
res.Failure().reason.Str(),
R"(:2:35 error: var: a builtin and location cannot be both declared for a module scope var
- %1:ptr<__in, vec4<f32>, read> = var @binding_point(0, 0) @location(0) @builtin(position)
+ %1:ptr<__in, vec4<f32>, read> = var @location(0) @builtin(position)
^^^
:1:1 note: in block
@@ -5174,7 +5221,7 @@
note: # Disassembly
$B1: { # root
- %1:ptr<__in, vec4<f32>, read> = var @binding_point(0, 0) @location(0) @builtin(position)
+ %1:ptr<__in, vec4<f32>, read> = var @location(0) @builtin(position)
}
)");
@@ -5190,7 +5237,6 @@
{mod.symbols.New("a"), ty.f32(), attr},
});
auto* v = b.Var(ty.ptr(AddressSpace::kOut, str_ty, read_write));
- v->SetBindingPoint(0, 0);
mod.root_block->Append(v);
auto res = ir::Validate(mod);
@@ -5198,7 +5244,7 @@
EXPECT_EQ(
res.Failure().reason.Str(),
R"(:6:41 error: var: a builtin and location cannot be both declared for a module scope var struct member
- %1:ptr<__out, MyStruct, read_write> = var @binding_point(0, 0)
+ %1:ptr<__out, MyStruct, read_write> = var
^^^
:5:1 note: in block
@@ -5211,7 +5257,7 @@
}
$B1: { # root
- %1:ptr<__out, MyStruct, read_write> = var @binding_point(0, 0)
+ %1:ptr<__out, MyStruct, read_write> = var
}
)");
diff --git a/src/tint/lang/glsl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/glsl/writer/raise/builtin_polyfill_test.cc
index 8a19095..c33a0c6 100644
--- a/src/tint/lang/glsl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/glsl/writer/raise/builtin_polyfill_test.cc
@@ -202,7 +202,6 @@
TEST_F(GlslWriter_BuiltinPolyfillTest, AtomicCompareExchangeWeak) {
auto* var = b.Var("v", workgroup, ty.atomic<i32>(), core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.ComputeFunction("foo");
@@ -219,7 +218,7 @@
}
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -239,7 +238,7 @@
}
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -261,7 +260,6 @@
TEST_F(GlslWriter_BuiltinPolyfillTest, AtomicSub) {
auto* var = b.Var("v", workgroup, ty.atomic<i32>(), core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.ComputeFunction("foo");
@@ -272,7 +270,7 @@
auto* src = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -287,7 +285,7 @@
auto* expect = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -306,7 +304,6 @@
TEST_F(GlslWriter_BuiltinPolyfillTest, AtomicSub_u32) {
auto* var = b.Var("v", workgroup, ty.atomic<u32>(), core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.ComputeFunction("foo");
@@ -317,7 +314,7 @@
auto* src = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<u32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<u32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -332,7 +329,7 @@
auto* expect = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<u32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<u32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -350,7 +347,6 @@
TEST_F(GlslWriter_BuiltinPolyfillTest, AtomicLoad) {
auto* var = b.Var("v", workgroup, ty.atomic<i32>(), core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.ComputeFunction("foo");
@@ -361,7 +357,7 @@
auto* src = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -376,7 +372,7 @@
auto* expect = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
diff --git a/src/tint/lang/glsl/writer/var_and_let_test.cc b/src/tint/lang/glsl/writer/var_and_let_test.cc
index bdf9288..c3605c3 100644
--- a/src/tint/lang/glsl/writer/var_and_let_test.cc
+++ b/src/tint/lang/glsl/writer/var_and_let_test.cc
@@ -123,7 +123,6 @@
core::IOAttributes attrs = {};
attrs.builtin = core::BuiltinValue::kLocalInvocationIndex;
v->SetAttributes(attrs);
- v->SetBindingPoint(1, 2);
});
ASSERT_TRUE(Generate()) << err_ << output_.glsl;
diff --git a/src/tint/lang/hlsl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/hlsl/writer/raise/builtin_polyfill_test.cc
index 1febf2b..65b2f52 100644
--- a/src/tint/lang/hlsl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/builtin_polyfill_test.cc
@@ -5775,7 +5775,6 @@
TEST_P(HlslBuiltinPolyfillWorkgroupAtomic, Access) {
auto param = GetParam();
auto* var = b.Var("v", workgroup, ty.atomic<i32>(), core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
@@ -5786,7 +5785,7 @@
std::string src = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @fragment func():void {
@@ -5802,7 +5801,7 @@
std::string expect = R"(
$B1: { # root
- %v:ptr<workgroup, atomic<i32>, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, atomic<i32>, read_write> = var
}
%foo = @fragment func():void {
@@ -5840,7 +5839,6 @@
});
auto* var = b.Var("v", workgroup, sb, core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
@@ -5858,7 +5856,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -5879,7 +5877,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -5903,7 +5901,6 @@
});
auto* var = b.Var("v", workgroup, sb, core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
@@ -5921,7 +5918,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -5943,7 +5940,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -5969,7 +5966,6 @@
});
auto* var = b.Var("v", workgroup, sb, core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
@@ -5989,7 +5985,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -6014,7 +6010,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -6047,7 +6043,6 @@
});
auto* var = b.Var("v", workgroup, sb, core::Access::kReadWrite);
- var->SetBindingPoint(0, 0);
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
@@ -6072,7 +6067,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
@@ -6099,7 +6094,7 @@
}
$B1: { # root
- %v:ptr<workgroup, SB, read_write> = var @binding_point(0, 0)
+ %v:ptr<workgroup, SB, read_write> = var
}
%foo = @fragment func():void {
diff --git a/src/tint/lang/hlsl/writer/var_let_test.cc b/src/tint/lang/hlsl/writer/var_let_test.cc
index a598027..ff18836 100644
--- a/src/tint/lang/hlsl/writer/var_let_test.cc
+++ b/src/tint/lang/hlsl/writer/var_let_test.cc
@@ -504,7 +504,6 @@
TEST_F(HlslWriterTest, VarPrivate) {
auto* s = b.Var("u", ty.ptr<private_>(ty.vec4<f32>()));
- s->SetBindingPoint(2, 1);
b.ir.root_block->Append(s);
@@ -520,7 +519,6 @@
TEST_F(HlslWriterTest, VarWorkgroup) {
auto* s = b.Var("u", ty.ptr<workgroup>(ty.vec4<f32>()));
- s->SetBindingPoint(2, 1);
b.ir.root_block->Append(s);
diff --git a/src/tint/lang/spirv/writer/var_test.cc b/src/tint/lang/spirv/writer/var_test.cc
index 3a9227a..2e8b159 100644
--- a/src/tint/lang/spirv/writer/var_test.cc
+++ b/src/tint/lang/spirv/writer/var_test.cc
@@ -304,7 +304,6 @@
TEST_F(SpirvWriterTest, StorageVar_Workgroup_WithVulkan) {
auto* v = b.Var("v", ty.ptr<workgroup, i32, read_write>());
- v->SetBindingPoint(0, 0);
mod.root_block->Append(v);
auto* func = b.ComputeFunction("foo");
