resolver: Validate pipline stage use for intrinsics Use the new [[stage()]] decorations in intrinsics.def to validate that intrinsics are only called from the correct pipeline stages. Fixed: tint:657 Change-Id: I9efda26369c45c6f816bdaa53408d3909db403a1 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53084 Reviewed-by: Ryan Harrison <rharrison@chromium.org> Reviewed-by: David Neto <dneto@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/resolver/builtins_validation_test.cc b/src/resolver/builtins_validation_test.cc index fe2b59d..d526f1e 100644 --- a/src/resolver/builtins_validation_test.cc +++ b/src/resolver/builtins_validation_test.cc
@@ -12,6 +12,7 @@ // See the License for the specific language governing permissions and // limitations under the License. +#include "src/ast/call_statement.h" #include "src/resolver/resolver_test_helper.h" namespace tint { @@ -283,7 +284,8 @@ params.push_back(Expr(1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->Is<sem::F32>()); @@ -298,7 +300,8 @@ params.push_back(vec2<f32>(1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector()); @@ -313,7 +316,8 @@ params.push_back(vec3<f32>(1.0f, 1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector()); @@ -328,7 +332,8 @@ params.push_back(vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc index 8c27acf..ba8af09 100644 --- a/src/resolver/intrinsic_test.cc +++ b/src/resolver/intrinsic_test.cc
@@ -55,7 +55,8 @@ Global("ident", ty.f32(), ast::StorageClass::kInput); auto* expr = Call(name, "ident"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); @@ -68,7 +69,8 @@ Global("ident", ty.vec4<f32>(), ast::StorageClass::kInput); auto* expr = Call(name, "ident"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); @@ -1927,7 +1929,8 @@ param.buildSamplerVariable(this); auto* call = Call(param.function, param.args(this)); - WrapInFunction(call); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); ASSERT_TRUE(r()->Resolve()) << r()->error();
diff --git a/src/resolver/intrinsic_validation_test.cc b/src/resolver/intrinsic_validation_test.cc new file mode 100644 index 0000000..31b5d3d --- /dev/null +++ b/src/resolver/intrinsic_validation_test.cc
@@ -0,0 +1,97 @@ +// Copyright 2021 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/resolver/resolver.h" + +#include "gmock/gmock.h" +#include "src/ast/assignment_statement.h" +#include "src/ast/bitcast_expression.h" +#include "src/ast/break_statement.h" +#include "src/ast/call_statement.h" +#include "src/ast/continue_statement.h" +#include "src/ast/if_statement.h" +#include "src/ast/intrinsic_texture_helper_test.h" +#include "src/ast/loop_statement.h" +#include "src/ast/return_statement.h" +#include "src/ast/stage_decoration.h" +#include "src/ast/struct_block_decoration.h" +#include "src/ast/switch_statement.h" +#include "src/ast/unary_op_expression.h" +#include "src/ast/variable_decl_statement.h" +#include "src/resolver/resolver_test_helper.h" +#include "src/sem/call.h" +#include "src/sem/function.h" +#include "src/sem/member_accessor_expression.h" +#include "src/sem/sampled_texture_type.h" +#include "src/sem/statement.h" +#include "src/sem/variable.h" + +using ::testing::ElementsAre; +using ::testing::HasSubstr; + +namespace tint { +namespace resolver { +namespace { + +using IntrinsicType = sem::IntrinsicType; + +using ResolverIntrinsicValidationTest = ResolverTest; + +TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageDirect) { + // [[stage(compute)]] fn func { return dpdx(1.0); } + + auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"), + ast::ExpressionList{Expr(1.0f)}); + Func(Source{{1, 2}}, "func", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(dpdx)}, + {Stage(ast::PipelineStage::kCompute)}); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), + "3:4 error: built-in cannot be used by compute pipeline stage"); +} + +TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageIndirect) { + // fn f0 { return dpdx(1.0); } + // fn f1 { f0(); } + // fn f2 { f1(); } + // [[stage(compute)]] fn main { return f2(); } + + auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"), + ast::ExpressionList{Expr(1.0f)}); + Func(Source{{1, 2}}, "f0", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(dpdx)}); + + Func(Source{{3, 4}}, "f1", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f0"))}); + + Func(Source{{5, 6}}, "f2", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f1"))}); + + Func(Source{{7, 8}}, "main", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f2"))}, + {Stage(ast::PipelineStage::kCompute)}); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), + R"(3:4 error: built-in cannot be used by compute pipeline stage +1:2 note: called by function 'f0' +3:4 note: called by function 'f1' +5:6 note: called by function 'f2' +7:8 note: called by entry point 'main')"); +} + +} // namespace +} // namespace resolver +} // namespace tint
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc index 2ac42f9..7b59942 100644 --- a/src/resolver/resolver.cc +++ b/src/resolver/resolver.cc
@@ -237,6 +237,10 @@ } } + if (!ValidatePipelineStages()) { + return false; + } + bool result = true; for (auto* node : builder_->ASTNodes().Objects()) { @@ -1129,6 +1133,10 @@ bool Resolver::Function(ast::Function* func) { auto* info = function_infos_.Create<FunctionInfo>(func); + if (func->IsEntryPoint()) { + entry_points_.emplace_back(info); + } + TINT_SCOPED_ASSIGNMENT(current_function_, info); variable_stack_.push_scope(); @@ -1707,6 +1715,10 @@ builder_->Sem().Add( call, builder_->create<sem::Call>(call, result, current_statement_)); SetType(call, result->ReturnType()); + + current_function_->intrinsic_calls.emplace_back( + IntrinsicCallInfo{call, result}); + return true; } @@ -2460,25 +2472,76 @@ expr_info_.emplace(expr, ExpressionInfo{type, type_name, current_statement_}); } +bool Resolver::ValidatePipelineStages() { + auto check_intrinsic_calls = [&](FunctionInfo* func, + FunctionInfo* entry_point) { + auto stage = entry_point->declaration->pipeline_stage(); + for (auto& call : func->intrinsic_calls) { + if (!call.intrinsic->SupportedStages().Contains(stage)) { + std::stringstream err; + err << "built-in cannot be used by " << stage << " pipeline stage"; + diagnostics_.add_error(err.str(), call.call->source()); + if (func != entry_point) { + TraverseCallChain(entry_point, func, [&](FunctionInfo* f) { + diagnostics_.add_note( + "called by function '" + + builder_->Symbols().NameFor(f->declaration->symbol()) + "'", + f->declaration->source()); + }); + diagnostics_.add_note("called by entry point '" + + builder_->Symbols().NameFor( + entry_point->declaration->symbol()) + + "'", + entry_point->declaration->source()); + } + return false; + } + } + return true; + }; + + for (auto* entry_point : entry_points_) { + if (!check_intrinsic_calls(entry_point, entry_point)) { + return false; + } + for (auto* func : entry_point->transitive_calls) { + if (!check_intrinsic_calls(func, entry_point)) { + return false; + } + } + } + return true; +} + +template <typename CALLBACK> +void Resolver::TraverseCallChain(FunctionInfo* from, + FunctionInfo* to, + CALLBACK&& callback) const { + for (auto* f : from->transitive_calls) { + if (f == to) { + callback(f); + return; + } + if (f->transitive_calls.contains(to)) { + TraverseCallChain(f, to, callback); + callback(f); + return; + } + } + TINT_ICE(diagnostics_) + << "TraverseCallChain() 'from' does not transitively call 'to'"; +} + void Resolver::CreateSemanticNodes() const { auto& sem = builder_->Sem(); // Collate all the 'ancestor_entry_points' - this is a map of function // symbol to all the entry points that transitively call the function. std::unordered_map<Symbol, std::vector<Symbol>> ancestor_entry_points; - for (auto* func : builder_->AST().Functions()) { - auto it = function_to_info_.find(func); - if (it == function_to_info_.end()) { - continue; // Resolver has likely errored. Process what we can. - } - - auto* info = it->second; - if (!func->IsEntryPoint()) { - continue; - } - for (auto* call : info->transitive_calls) { + for (auto* entry_point : entry_points_) { + for (auto* call : entry_point->transitive_calls) { auto& vec = ancestor_entry_points[call->declaration->symbol()]; - vec.emplace_back(func->symbol()); + vec.emplace_back(entry_point->declaration->symbol()); } }
diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h index 252f8bc..2efdf7c 100644 --- a/src/resolver/resolver.h +++ b/src/resolver/resolver.h
@@ -51,6 +51,7 @@ } // namespace ast namespace sem { class Array; +class Intrinsic; class Statement; } // namespace sem @@ -100,6 +101,11 @@ sem::BindingPoint binding_point; }; + struct IntrinsicCallInfo { + const ast::CallExpression* call; + const sem::Intrinsic* intrinsic; + }; + /// Structure holding semantic information about a function. /// Used to build the sem::Function nodes at the end of resolving. struct FunctionInfo { @@ -115,9 +121,13 @@ sem::Type* return_type = nullptr; std::string return_type_name; std::array<sem::WorkgroupDimension, 3> workgroup_size; + std::vector<IntrinsicCallInfo> intrinsic_calls; // List of transitive calls this function makes UniqueVector<FunctionInfo*> transitive_calls; + + // List of entry point functions that transitively call this function + UniqueVector<FunctionInfo*> ancestor_entry_points; }; /// Structure holding semantic information about an expression. @@ -183,6 +193,8 @@ /// @returns true on success, false on error bool ResolveInternal(); + bool ValidatePipelineStages(); + /// Creates the nodes and adds them to the sem::Info mappings of the /// ProgramBuilder. void CreateSemanticNodes() const; @@ -359,12 +371,18 @@ /// @param node the AST node. void Mark(const ast::Node* node); + template <typename CALLBACK> + void TraverseCallChain(FunctionInfo* from, + FunctionInfo* to, + CALLBACK&& callback) const; + ProgramBuilder* const builder_; diag::List& diagnostics_; std::unique_ptr<IntrinsicTable> const intrinsic_table_; sem::BlockStatement* current_block_ = nullptr; ScopeStack<VariableInfo*> variable_stack_; std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_; + std::vector<FunctionInfo*> entry_points_; std::unordered_map<const ast::Function*, FunctionInfo*> function_to_info_; std::unordered_map<const ast::Variable*, VariableInfo*> variable_to_info_; std::unordered_map<ast::CallExpression*, FunctionCallInfo> function_calls_;