encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Add semantic::Function, use it. 

Pull the mutable semantic fields from ast::Function and into a new semantic::Function node.
Have the TypeDeterminer create these semantic::Function nodes.

Bug: tint:390
Change-Id: I237b1bed8709dd9a3cfa24d85d48fc77b7e532da
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/39902
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton 2021-02-03 16:43:20 +00:00 committed by Commit Bot service account
parent c694d43c75
commit 87c78ddabc
22 changed files with 744 additions and 512 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
BUILD.gn
View File

@ -382,6 +382,7 @@ source_set("libtint_core_src") {
"src/semantic/info.h", "src/semantic/info.h",
"src/semantic/node.h", "src/semantic/node.h",
"src/semantic/sem_expression.cc", "src/semantic/sem_expression.cc",
"src/semantic/sem_function.cc",
"src/semantic/sem_info.cc", "src/semantic/sem_info.cc",
"src/semantic/sem_node.cc", "src/semantic/sem_node.cc",
"src/semantic/type_mappings.h", "src/semantic/type_mappings.h",

1
src/CMakeLists.txt
View File

@ -196,6 +196,7 @@ set(TINT_LIB_SRCS
semantic/info.h semantic/info.h
semantic/node.h semantic/node.h
semantic/sem_expression.cc semantic/sem_expression.cc
semantic/sem_function.cc
semantic/sem_info.cc semantic/sem_info.cc
semantic/sem_node.cc semantic/sem_node.cc
semantic/type_mappings.h semantic/type_mappings.h

222
src/ast/function.cc
View File

@ -65,161 +65,6 @@ PipelineStage Function::pipeline_stage() const {
return PipelineStage::kNone; return PipelineStage::kNone;
} }
void Function::add_referenced_module_variable(Variable* var) {
for (const auto* v : referenced_module_vars_) {
if (v->symbol() == var->symbol()) {
return;
}
}
referenced_module_vars_.push_back(var);
}
void Function::add_local_referenced_module_variable(Variable* var) {
for (const auto* v : local_referenced_module_vars_) {
if (v->symbol() == var->symbol()) {
return;
}
}
local_referenced_module_vars_.push_back(var);
}
const std::vector<std::pair<Variable*, LocationDecoration*>>
Function::referenced_location_variables() const {
std::vector<std::pair<Variable*, LocationDecoration*>> ret;
for (auto* var : referenced_module_variables()) {
for (auto* deco : var->decorations()) {
if (auto* location = deco->As<LocationDecoration>()) {
ret.push_back({var, location});
break;
}
}
}
return ret;
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_uniform_variables() const {
std::vector<std::pair<Variable*, Function::BindingInfo>> ret;
for (auto* var : referenced_module_variables()) {
if (var->storage_class() != StorageClass::kUniform) {
continue;
}
BindingDecoration* binding = nullptr;
GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<BindingDecoration>()) {
binding = b;
} else if (auto* g = deco->As<GroupDecoration>()) {
group = g;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_storagebuffer_variables() const {
std::vector<std::pair<Variable*, Function::BindingInfo>> ret;
for (auto* var : referenced_module_variables()) {
if (var->storage_class() != StorageClass::kStorage) {
continue;
}
BindingDecoration* binding = nullptr;
GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<Variable*, BuiltinDecoration*>>
Function::referenced_builtin_variables() const {
std::vector<std::pair<Variable*, BuiltinDecoration*>> ret;
for (auto* var : referenced_module_variables()) {
for (auto* deco : var->decorations()) {
if (auto* builtin = deco->As<BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
}
}
}
return ret;
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_sampler_variables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kSampler);
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_comparison_sampler_variables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kComparisonSampler);
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_sampled_texture_variables() const {
return ReferencedSampledTextureVariablesImpl(false);
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::referenced_multisampled_texture_variables() const {
return ReferencedSampledTextureVariablesImpl(true);
}
const std::vector<std::pair<Variable*, BuiltinDecoration*>>
Function::local_referenced_builtin_variables() const {
std::vector<std::pair<Variable*, BuiltinDecoration*>> ret;
for (auto* var : local_referenced_module_variables()) {
for (auto* deco : var->decorations()) {
if (auto* builtin = deco->As<BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
}
}
}
return ret;
}
void Function::add_ancestor_entry_point(Symbol ep) {
for (const auto& point : ancestor_entry_points_) {
if (point == ep) {
return;
}
}
ancestor_entry_points_.push_back(ep);
}
bool Function::HasAncestorEntryPoint(Symbol symbol) const {
for (const auto& point : ancestor_entry_points_) {
if (point == symbol) {
return true;
}
}
return false;
}
const Statement* Function::get_last_statement() const { const Statement* Function::get_last_statement() const {
return body_->last(); return body_->last();
} }
@ -295,73 +140,6 @@ std::string Function::type_name() const {
return out.str(); return out.str();
} }
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
std::vector<std::pair<Variable*, Function::BindingInfo>> ret;
for (auto* var : referenced_module_variables()) {
auto* unwrapped_type = var->type()->UnwrapIfNeeded();
auto* sampler = unwrapped_type->As<type::Sampler>();
if (sampler == nullptr || sampler->kind() != kind) {
continue;
}
BindingDecoration* binding = nullptr;
GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<BindingDecoration>()) {
binding = b;
}
if (auto* s = deco->As<GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<Variable*, Function::BindingInfo>>
Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const {
std::vector<std::pair<Variable*, Function::BindingInfo>> ret;
for (auto* var : referenced_module_variables()) {
auto* unwrapped_type = var->type()->UnwrapIfNeeded();
auto* texture = unwrapped_type->As<type::Texture>();
if (texture == nullptr) {
continue;
}
auto is_multisampled = texture->Is<type::MultisampledTexture>();
auto is_sampled = texture->Is<type::SampledTexture>();
if ((multisampled && !is_multisampled) || (!multisampled && !is_sampled)) {
continue;
}
BindingDecoration* binding = nullptr;
GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
Function* FunctionList::Find(Symbol sym) const { Function* FunctionList::Find(Symbol sym) const {
for (auto* func : *this) { for (auto* func : *this) {
if (func->symbol() == sym) { if (func->symbol() == sym) {

93
src/ast/function.h
View File

@ -43,14 +43,6 @@ namespace ast {
/// A Function statement. /// A Function statement.
class Function : public Castable<Function, Node> { class Function : public Castable<Function, Node> {
public: public:
/// Information about a binding
struct BindingInfo {
/// The binding decoration
BindingDecoration* binding = nullptr;
/// The group decoration
GroupDecoration* group = nullptr;
};
/// Create a function /// Create a function
/// @param source the variable source /// @param source the variable source
/// @param symbol the function symbol /// @param symbol the function symbol
@ -87,82 +79,9 @@ class Function : public Castable<Function, Node> {
/// @returns true if this function is an entry point /// @returns true if this function is an entry point
bool IsEntryPoint() const { return pipeline_stage() != PipelineStage::kNone; } bool IsEntryPoint() const { return pipeline_stage() != PipelineStage::kNone; }
/// Adds the given variable to the list of referenced module variables if it
/// is not already included.
/// @param var the module variable to add
void add_referenced_module_variable(Variable* var);
/// Adds the given variable to the list of locally referenced module variables
/// if it is not already included.
/// @param var the module variable to add
void add_local_referenced_module_variable(Variable* var);
/// Note: If this function calls other functions, the return will also include
/// all of the referenced variables from the callees.
/// @returns the referenced module variables
const std::vector<Variable*>& referenced_module_variables() const {
return referenced_module_vars_;
}
/// @returns the locally referenced module variables
const std::vector<Variable*>& local_referenced_module_variables() const {
return local_referenced_module_vars_;
}
/// Retrieves any referenced location variables
/// @returns the <variable, decoration> pair.
const std::vector<std::pair<Variable*, LocationDecoration*>>
referenced_location_variables() const;
/// Retrieves any referenced builtin variables
/// @returns the <variable, decoration> pair.
const std::vector<std::pair<Variable*, BuiltinDecoration*>>
referenced_builtin_variables() const;
/// Retrieves any referenced uniform variables. Note, the variables must be
/// decorated with both binding and group decorations.
/// @returns the referenced uniforms
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_uniform_variables() const;
/// Retrieves any referenced storagebuffer variables. Note, the variables
/// must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_storagebuffer_variables() const;
/// Retrieves any referenced regular Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_sampler_variables() const;
/// Retrieves any referenced comparison Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_comparison_sampler_variables() const;
/// Retrieves any referenced sampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_sampled_texture_variables() const;
/// Retrieves any referenced multisampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
const std::vector<std::pair<Variable*, Function::BindingInfo>>
referenced_multisampled_texture_variables() const;
/// Retrieves any locally referenced builtin variables
/// @returns the <variable, decoration> pairs.
const std::vector<std::pair<Variable*, BuiltinDecoration*>>
local_referenced_builtin_variables() const;
/// Adds an ancestor entry point
/// @param ep the entry point ancestor
void add_ancestor_entry_point(Symbol ep);
/// @returns the ancestor entry points
const std::vector<Symbol>& ancestor_entry_points() const {
return ancestor_entry_points_;
}
/// Checks if the given entry point is an ancestor
/// @param sym the entry point symbol
/// @returns true if `sym` is an ancestor entry point of this function
bool HasAncestorEntryPoint(Symbol sym) const;
/// @returns the function return type. /// @returns the function return type.
type::Type* return_type() const { return return_type_; } type::Type* return_type() const { return return_type_; }
/// @returns a pointer to the last statement of the function or nullptr if /// @returns a pointer to the last statement of the function or nullptr if
// function is empty // function is empty
const Statement* get_last_statement() const; const Statement* get_last_statement() const;
@ -196,20 +115,12 @@ class Function : public Castable<Function, Node> {
private: private:
Function(const Function&) = delete; Function(const Function&) = delete;
const std::vector<std::pair<Variable*, Function::BindingInfo>>
ReferencedSamplerVariablesImpl(type::SamplerKind kind) const;
const std::vector<std::pair<Variable*, Function::BindingInfo>>
ReferencedSampledTextureVariablesImpl(bool multisampled) const;
Symbol const symbol_; Symbol const symbol_;
VariableList const params_; VariableList const params_;
type::Type* const return_type_; type::Type* const return_type_;
BlockStatement* const body_; BlockStatement* const body_;
FunctionDecorationList const decorations_;
std::vector<Variable*> referenced_module_vars_; // Semantic info
std::vector<Variable*> local_referenced_module_vars_; // Semantic info
std::vector<Symbol> ancestor_entry_points_; // Semantic info
FunctionDecorationList decorations_; // Semantic info
}; };
/// A list of functions /// A list of functions

108
src/ast/function_test.cc
View File

@ -53,114 +53,6 @@ TEST_F(FunctionTest, Creation_WithSource) {
EXPECT_EQ(src.range.begin.column, 2u); EXPECT_EQ(src.range.begin.column, 2u);
} }
TEST_F(FunctionTest, AddDuplicateReferencedVariables) {
auto* v = Var("var", StorageClass::kInput, ty.i32());
auto* f = Func("func", VariableList{}, ty.void_(), StatementList{},
FunctionDecorationList{});
f->add_referenced_module_variable(v);
ASSERT_EQ(f->referenced_module_variables().size(), 1u);
EXPECT_EQ(f->referenced_module_variables()[0], v);
f->add_referenced_module_variable(v);
ASSERT_EQ(f->referenced_module_variables().size(), 1u);
auto* v2 = Var("var2", StorageClass::kOutput, ty.i32());
f->add_referenced_module_variable(v2);
ASSERT_EQ(f->referenced_module_variables().size(), 2u);
EXPECT_EQ(f->referenced_module_variables()[1], v2);
}
TEST_F(FunctionTest, GetReferenceLocations) {
auto* loc1 = Var("loc1", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<LocationDecoration>(0),
});
auto* loc2 = Var("loc2", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<LocationDecoration>(1),
});
auto* builtin1 = Var("builtin1", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<BuiltinDecoration>(Builtin::kPosition),
});
auto* builtin2 = Var("builtin2", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<BuiltinDecoration>(Builtin::kFragDepth),
});
auto* f = Func("func", VariableList{}, ty.void_(), StatementList{},
FunctionDecorationList{});
f->add_referenced_module_variable(loc1);
f->add_referenced_module_variable(builtin1);
f->add_referenced_module_variable(loc2);
f->add_referenced_module_variable(builtin2);
ASSERT_EQ(f->referenced_module_variables().size(), 4u);
auto ref_locs = f->referenced_location_variables();
ASSERT_EQ(ref_locs.size(), 2u);
EXPECT_EQ(ref_locs[0].first, loc1);
EXPECT_EQ(ref_locs[0].second->value(), 0u);
EXPECT_EQ(ref_locs[1].first, loc2);
EXPECT_EQ(ref_locs[1].second->value(), 1u);
}
TEST_F(FunctionTest, GetReferenceBuiltins) {
auto* loc1 = Var("loc1", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<LocationDecoration>(0),
});
auto* loc2 = Var("loc2", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<LocationDecoration>(1),
});
auto* builtin1 = Var("builtin1", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<BuiltinDecoration>(Builtin::kPosition),
});
auto* builtin2 = Var("builtin2", StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<BuiltinDecoration>(Builtin::kFragDepth),
});
auto* f = Func("func", VariableList{}, ty.void_(), StatementList{},
FunctionDecorationList{});
f->add_referenced_module_variable(loc1);
f->add_referenced_module_variable(builtin1);
f->add_referenced_module_variable(loc2);
f->add_referenced_module_variable(builtin2);
ASSERT_EQ(f->referenced_module_variables().size(), 4u);
auto ref_locs = f->referenced_builtin_variables();
ASSERT_EQ(ref_locs.size(), 2u);
EXPECT_EQ(ref_locs[0].first, builtin1);
EXPECT_EQ(ref_locs[0].second->value(), Builtin::kPosition);
EXPECT_EQ(ref_locs[1].first, builtin2);
EXPECT_EQ(ref_locs[1].second->value(), Builtin::kFragDepth);
}
TEST_F(FunctionTest, AddDuplicateEntryPoints) {
auto* f = Func("func", VariableList{}, ty.void_(), StatementList{},
FunctionDecorationList{});
auto main_sym = Symbols().Get("main");
f->add_ancestor_entry_point(main_sym);
ASSERT_EQ(1u, f->ancestor_entry_points().size());
EXPECT_EQ(main_sym, f->ancestor_entry_points()[0]);
f->add_ancestor_entry_point(main_sym);
ASSERT_EQ(1u, f->ancestor_entry_points().size());
EXPECT_EQ(main_sym, f->ancestor_entry_points()[0]);
}
TEST_F(FunctionTest, IsValid) { TEST_F(FunctionTest, IsValid) {
VariableList params; VariableList params;
params.push_back(Var("var", StorageClass::kNone, ty.i32())); params.push_back(Var("var", StorageClass::kNone, ty.i32()));

30
src/inspector/inspector.cc
View File

@ -29,6 +29,7 @@
#include "src/ast/uint_literal.h" #include "src/ast/uint_literal.h"
#include "src/ast/variable.h" #include "src/ast/variable.h"
#include "src/program.h" #include "src/program.h"
#include "src/semantic/function.h"
#include "src/type/access_control_type.h" #include "src/type/access_control_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
#include "src/type/f32_type.h" #include "src/type/f32_type.h"
@ -64,7 +65,7 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
std::tie(entry_point.workgroup_size_x, entry_point.workgroup_size_y, std::tie(entry_point.workgroup_size_x, entry_point.workgroup_size_y,
entry_point.workgroup_size_z) = func->workgroup_size(); entry_point.workgroup_size_z) = func->workgroup_size();
for (auto* var : func->referenced_module_variables()) { for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) {
auto name = program_->Symbols().NameFor(var->symbol()); auto name = program_->Symbols().NameFor(var->symbol());
if (var->HasBuiltinDecoration()) { if (var->HasBuiltinDecoration()) {
continue; continue;
@ -185,10 +186,11 @@ std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
std::vector<ResourceBinding> result; std::vector<ResourceBinding> result;
for (auto& ruv : func->referenced_uniform_variables()) { auto* func_sem = program_->Sem().Get(func);
for (auto& ruv : func_sem->ReferencedUniformVariables()) {
ResourceBinding entry; ResourceBinding entry;
ast::Variable* var = nullptr; ast::Variable* var = nullptr;
ast::Function::BindingInfo binding_info; semantic::Function::BindingInfo binding_info;
std::tie(var, binding_info) = ruv; std::tie(var, binding_info) = ruv;
if (!var->type()->Is<type::AccessControl>()) { if (!var->type()->Is<type::AccessControl>()) {
continue; continue;
@ -235,10 +237,11 @@ std::vector<ResourceBinding> Inspector::GetSamplerResourceBindings(
std::vector<ResourceBinding> result; std::vector<ResourceBinding> result;
for (auto& rs : func->referenced_sampler_variables()) { auto* func_sem = program_->Sem().Get(func);
for (auto& rs : func_sem->ReferencedSamplerVariables()) {
ResourceBinding entry; ResourceBinding entry;
ast::Variable* var = nullptr; ast::Variable* var = nullptr;
ast::Function::BindingInfo binding_info; semantic::Function::BindingInfo binding_info;
std::tie(var, binding_info) = rs; std::tie(var, binding_info) = rs;
entry.bind_group = binding_info.group->value(); entry.bind_group = binding_info.group->value();
@ -259,10 +262,11 @@ std::vector<ResourceBinding> Inspector::GetComparisonSamplerResourceBindings(
std::vector<ResourceBinding> result; std::vector<ResourceBinding> result;
for (auto& rcs : func->referenced_comparison_sampler_variables()) { auto* func_sem = program_->Sem().Get(func);
for (auto& rcs : func_sem->ReferencedComparisonSamplerVariables()) {
ResourceBinding entry; ResourceBinding entry;
ast::Variable* var = nullptr; ast::Variable* var = nullptr;
ast::Function::BindingInfo binding_info; semantic::Function::BindingInfo binding_info;
std::tie(var, binding_info) = rcs; std::tie(var, binding_info) = rcs;
entry.bind_group = binding_info.group->value(); entry.bind_group = binding_info.group->value();
@ -307,11 +311,12 @@ std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
return {}; return {};
} }
auto* func_sem = program_->Sem().Get(func);
std::vector<ResourceBinding> result; std::vector<ResourceBinding> result;
for (auto& rsv : func->referenced_storagebuffer_variables()) { for (auto& rsv : func_sem->ReferencedStoragebufferVariables()) {
ResourceBinding entry; ResourceBinding entry;
ast::Variable* var = nullptr; ast::Variable* var = nullptr;
ast::Function::BindingInfo binding_info; semantic::Function::BindingInfo binding_info;
std::tie(var, binding_info) = rsv; std::tie(var, binding_info) = rsv;
auto* ac_type = var->type()->As<type::AccessControl>(); auto* ac_type = var->type()->As<type::AccessControl>();
@ -347,13 +352,14 @@ std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
} }
std::vector<ResourceBinding> result; std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
auto& referenced_variables = auto& referenced_variables =
multisampled_only ? func->referenced_multisampled_texture_variables() multisampled_only ? func_sem->ReferencedMultisampledTextureVariables()
: func->referenced_sampled_texture_variables(); : func_sem->ReferencedSampledTextureVariables();
for (auto& ref : referenced_variables) { for (auto& ref : referenced_variables) {
ResourceBinding entry; ResourceBinding entry;
ast::Variable* var = nullptr; ast::Variable* var = nullptr;
ast::Function::BindingInfo binding_info; semantic::Function::BindingInfo binding_info;
std::tie(var, binding_info) = ref; std::tie(var, binding_info) = ref;
entry.bind_group = binding_info.group->value(); entry.bind_group = binding_info.group->value();

2
src/semantic/expression.h
View File

@ -17,8 +17,6 @@
#include "src/semantic/node.h" #include "src/semantic/node.h"
#include "src/semantic/type_mappings.h"
namespace tint { namespace tint {
// Forward declarations // Forward declarations

148
src/semantic/function.h Normal file
View File

@ -0,0 +1,148 @@
// 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.
#ifndef SRC_SEMANTIC_FUNCTION_H_
#define SRC_SEMANTIC_FUNCTION_H_
#include <utility>
#include <vector>
#include "src/semantic/node.h"
#include "src/type/sampler_type.h"
namespace tint {
// Forward declarations
namespace ast {
class BindingDecoration;
class GroupDecoration;
class Variable;
class LocationDecoration;
class BuiltinDecoration;
} // namespace ast
namespace type {
class Type;
} // namespace type
namespace semantic {
/// Function holds the semantic information for function nodes.
class Function : public Castable<Function, Node> {
public:
/// Information about a binding
struct BindingInfo {
/// The binding decoration
ast::BindingDecoration* binding = nullptr;
/// The group decoration
ast::GroupDecoration* group = nullptr;
};
/// Constructor
/// @param referenced_module_vars the referenced module variables
/// @param local_referenced_module_vars the locally referenced module
/// variables
/// @param ancestor_entry_points the ancestor entry points
explicit Function(std::vector<ast::Variable*> referenced_module_vars,
std::vector<ast::Variable*> local_referenced_module_vars,
std::vector<Symbol> ancestor_entry_points);
/// Destructor
~Function() override;
/// Note: If this function calls other functions, the return will also include
/// all of the referenced variables from the callees.
/// @returns the referenced module variables
const std::vector<ast::Variable*>& ReferencedModuleVariables() const {
return referenced_module_vars_;
}
/// @returns the locally referenced module variables
const std::vector<ast::Variable*>& LocalReferencedModuleVariables() const {
return local_referenced_module_vars_;
}
/// @returns the ancestor entry points
const std::vector<Symbol>& AncestorEntryPoints() const {
return ancestor_entry_points_;
}
/// Retrieves any referenced location variables
/// @returns the <variable, decoration> pair.
const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
ReferencedLocationVariables() const;
/// Retrieves any referenced builtin variables
/// @returns the <variable, decoration> pair.
const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
ReferencedBuiltinVariables() const;
/// Retrieves any referenced uniform variables. Note, the variables must be
/// decorated with both binding and group decorations.
/// @returns the referenced uniforms
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedUniformVariables() const;
/// Retrieves any referenced storagebuffer variables. Note, the variables
/// must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedStoragebufferVariables() const;
/// Retrieves any referenced regular Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedSamplerVariables() const;
/// Retrieves any referenced comparison Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedComparisonSamplerVariables() const;
/// Retrieves any referenced sampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedSampledTextureVariables() const;
/// Retrieves any referenced multisampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedMultisampledTextureVariables() const;
/// Retrieves any locally referenced builtin variables
/// @returns the <variable, decoration> pairs.
const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
LocalReferencedBuiltinVariables() const;
/// Checks if the given entry point is an ancestor
/// @param sym the entry point symbol
/// @returns true if `sym` is an ancestor entry point of this function
bool HasAncestorEntryPoint(Symbol sym) const;
private:
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedSamplerVariablesImpl(type::SamplerKind kind) const;
const std::vector<std::pair<ast::Variable*, BindingInfo>>
ReferencedSampledTextureVariablesImpl(bool multisampled) const;
std::vector<ast::Variable*> const referenced_module_vars_;
std::vector<ast::Variable*> const local_referenced_module_vars_;
std::vector<Symbol> const ancestor_entry_points_;
};
} // namespace semantic
} // namespace tint
#endif // SRC_SEMANTIC_FUNCTION_H_

237
src/semantic/sem_function.cc Normal file
View File

@ -0,0 +1,237 @@
// 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/semantic/function.h"
#include "src/ast/binding_decoration.h"
#include "src/ast/builtin_decoration.h"
#include "src/ast/group_decoration.h"
#include "src/ast/location_decoration.h"
#include "src/ast/variable.h"
#include "src/ast/variable_decoration.h"
#include "src/type/multisampled_texture_type.h"
#include "src/type/sampled_texture_type.h"
#include "src/type/texture_type.h"
TINT_INSTANTIATE_CLASS_ID(tint::semantic::Function);
namespace tint {
namespace semantic {
Function::Function(std::vector<ast::Variable*> referenced_module_vars,
std::vector<ast::Variable*> local_referenced_module_vars,
std::vector<Symbol> ancestor_entry_points)
: referenced_module_vars_(std::move(referenced_module_vars)),
local_referenced_module_vars_(std::move(local_referenced_module_vars)),
ancestor_entry_points_(std::move(ancestor_entry_points)) {}
Function::~Function() = default;
const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
Function::ReferencedLocationVariables() const {
std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* deco : var->decorations()) {
if (auto* location = deco->As<ast::LocationDecoration>()) {
ret.push_back({var, location});
break;
}
}
}
return ret;
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedUniformVariables() const {
std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
if (var->storage_class() != ast::StorageClass::kUniform) {
continue;
}
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
} else if (auto* g = deco->As<ast::GroupDecoration>()) {
group = g;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedStoragebufferVariables() const {
std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
if (var->storage_class() != ast::StorageClass::kStorage) {
continue;
}
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<ast::GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
Function::ReferencedBuiltinVariables() const {
std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* deco : var->decorations()) {
if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
}
}
}
return ret;
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedSamplerVariables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kSampler);
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedComparisonSamplerVariables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kComparisonSampler);
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedSampledTextureVariables() const {
return ReferencedSampledTextureVariablesImpl(false);
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedMultisampledTextureVariables() const {
return ReferencedSampledTextureVariablesImpl(true);
}
const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
Function::LocalReferencedBuiltinVariables() const {
std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
for (auto* var : LocalReferencedModuleVariables()) {
for (auto* deco : var->decorations()) {
if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
}
}
}
return ret;
}
bool Function::HasAncestorEntryPoint(Symbol symbol) const {
for (const auto& point : ancestor_entry_points_) {
if (point == symbol) {
return true;
}
}
return false;
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
auto* unwrapped_type = var->type()->UnwrapIfNeeded();
auto* sampler = unwrapped_type->As<type::Sampler>();
if (sampler == nullptr || sampler->kind() != kind) {
continue;
}
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
}
if (auto* s = deco->As<ast::GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const {
std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
auto* unwrapped_type = var->type()->UnwrapIfNeeded();
auto* texture = unwrapped_type->As<type::Texture>();
if (texture == nullptr) {
continue;
}
auto is_multisampled = texture->Is<type::MultisampledTexture>();
auto is_sampled = texture->Is<type::SampledTexture>();
if ((multisampled && !is_multisampled) || (!multisampled && !is_sampled)) {
continue;
}
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
for (auto* deco : var->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<ast::GroupDecoration>()) {
group = s;
}
}
if (binding == nullptr || group == nullptr) {
continue;
}
ret.push_back({var, BindingInfo{binding, group}});
}
return ret;
}
} // namespace semantic
} // namespace tint

107
src/semantic/sem_function_test.cc Normal file
View File

@ -0,0 +1,107 @@
// 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/semantic/function.h"
#include "src/ast/builtin_decoration.h"
#include "src/ast/location_decoration.h"
#include "src/semantic/test_helper.h"
namespace tint {
namespace semantic {
namespace {
using FunctionTest = TestHelper;
TEST_F(FunctionTest, GetReferenceLocations) {
auto* loc1 = Var("loc1", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::LocationDecoration>(0),
});
auto* loc2 = Var("loc2", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::LocationDecoration>(1),
});
auto* builtin1 =
Var("builtin1", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::BuiltinDecoration>(ast::Builtin::kPosition),
});
auto* builtin2 =
Var("builtin2", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
});
auto* f = create<Function>(
/* referenced_module_vars */ std::vector<ast::Variable*>{loc1, builtin1,
loc2, builtin2},
/* local_referenced_module_vars */ std::vector<ast::Variable*>{},
/* ancestor_entry_points */ std::vector<Symbol>{});
ASSERT_EQ(f->ReferencedModuleVariables().size(), 4u);
auto ref_locs = f->ReferencedLocationVariables();
ASSERT_EQ(ref_locs.size(), 2u);
EXPECT_EQ(ref_locs[0].first, loc1);
EXPECT_EQ(ref_locs[0].second->value(), 0u);
EXPECT_EQ(ref_locs[1].first, loc2);
EXPECT_EQ(ref_locs[1].second->value(), 1u);
}
TEST_F(FunctionTest, GetReferenceBuiltins) {
auto* loc1 = Var("loc1", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::LocationDecoration>(0),
});
auto* loc2 = Var("loc2", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::LocationDecoration>(1),
});
auto* builtin1 =
Var("builtin1", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::BuiltinDecoration>(ast::Builtin::kPosition),
});
auto* builtin2 =
Var("builtin2", ast::StorageClass::kInput, ty.i32(), nullptr,
ast::VariableDecorationList{
create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
});
auto* f = create<Function>(
/* referenced_module_vars */ std::vector<ast::Variable*>{loc1, builtin1,
loc2, builtin2},
/* local_referenced_module_vars */ std::vector<ast::Variable*>{},
/* ancestor_entry_points */ std::vector<Symbol>{});
ASSERT_EQ(f->ReferencedModuleVariables().size(), 4u);
auto ref_locs = f->ReferencedBuiltinVariables();
ASSERT_EQ(ref_locs.size(), 2u);
EXPECT_EQ(ref_locs[0].first, builtin1);
EXPECT_EQ(ref_locs[0].second->value(), ast::Builtin::kPosition);
EXPECT_EQ(ref_locs[1].first, builtin2);
EXPECT_EQ(ref_locs[1].second->value(), ast::Builtin::kFragDepth);
}
} // namespace
} // namespace semantic
} // namespace tint

39
src/semantic/test_helper.h Normal file
View File

@ -0,0 +1,39 @@
// Copyright 2020 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.
#ifndef SRC_SEMANTIC_TEST_HELPER_H_
#define SRC_SEMANTIC_TEST_HELPER_H_
#include <memory>
#include <string>
#include <utility>
#include "gtest/gtest.h"
#include "src/program_builder.h"
namespace tint {
namespace semantic {
/// Helper class for testing
template <typename BASE>
class TestHelperBase : public BASE, public ProgramBuilder {};
using TestHelper = TestHelperBase<testing::Test>;
template <typename T>
using TestParamHelper = TestHelperBase<testing::TestWithParam<T>>;
} // namespace semantic
} // namespace tint
#endif // SRC_SEMANTIC_TEST_HELPER_H_

3
src/semantic/type_mappings.h
View File

@ -23,12 +23,14 @@ namespace tint {
namespace ast { namespace ast {
class Expression; class Expression;
class Function;
} // namespace ast } // namespace ast
namespace semantic { namespace semantic {
class Expression; class Expression;
class Function;
/// TypeMappings is a struct that holds dummy `operator()` methods that's used /// TypeMappings is a struct that holds dummy `operator()` methods that's used
/// by SemanticNodeTypeFor to map AST node types to their corresponding semantic /// by SemanticNodeTypeFor to map AST node types to their corresponding semantic
@ -38,6 +40,7 @@ class Expression;
struct TypeMappings { struct TypeMappings {
//! @cond Doxygen_Suppress //! @cond Doxygen_Suppress
semantic::Expression* operator()(ast::Expression*); semantic::Expression* operator()(ast::Expression*);
semantic::Function* operator()(ast::Function*);
//! @endcond //! @endcond
}; };

4
src/transform/first_index_offset.cc
View File

@ -49,6 +49,7 @@
#include "src/clone_context.h" #include "src/clone_context.h"
#include "src/program.h" #include "src/program.h"
#include "src/program_builder.h" #include "src/program_builder.h"
#include "src/semantic/function.h"
#include "src/type/struct_type.h" #include "src/type/struct_type.h"
#include "src/type/u32_type.h" #include "src/type/u32_type.h"
#include "src/type_determiner.h" #include "src/type_determiner.h"
@ -143,9 +144,10 @@ Transform::Output FirstIndexOffset::Run(const Program* in) {
if (buffer_var == nullptr) { if (buffer_var == nullptr) {
return nullptr; // no transform need, just clone func return nullptr; // no transform need, just clone func
} }
auto* func_sem = in->Sem().Get(func);
ast::StatementList statements; ast::StatementList statements;
for (const auto& data : for (const auto& data :
func->local_referenced_builtin_variables()) { func_sem->LocalReferencedBuiltinVariables()) {
if (data.second->value() == ast::Builtin::kVertexIndex) { if (data.second->value() == ast::Builtin::kVertexIndex) {
statements.emplace_back(CreateFirstIndexOffset( statements.emplace_back(CreateFirstIndexOffset(
in->Symbols().NameFor(vertex_index_sym), kFirstVertexName, in->Symbols().NameFor(vertex_index_sym), kFirstVertexName,

48
src/type_determiner.cc
View File

@ -44,6 +44,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/program_builder.h" #include "src/program_builder.h"
#include "src/semantic/expression.h" #include "src/semantic/expression.h"
#include "src/semantic/function.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
#include "src/type/bool_type.h" #include "src/type/bool_type.h"
#include "src/type/depth_texture_type.h" #include "src/type/depth_texture_type.h"
@ -97,9 +98,9 @@ void TypeDeterminer::set_referenced_from_function_if_needed(ast::Variable* var,
return; return;
} }
current_function_->add_referenced_module_variable(var); current_function_->referenced_module_vars.Add(var);
if (local) { if (local) {
current_function_->add_local_referenced_module_variable(var); current_function_->local_referenced_module_vars.Add(var);
} }
} }
@ -145,11 +146,14 @@ bool TypeDeterminer::Determine() {
} }
} }
CreateSemanticFunctions();
return true; return true;
} }
void TypeDeterminer::set_entry_points(const Symbol& fn_sym, Symbol ep_sym) { void TypeDeterminer::set_entry_points(const Symbol& fn_sym, Symbol ep_sym) {
symbol_to_function_[fn_sym]->add_ancestor_entry_point(ep_sym); auto* info = symbol_to_function_.at(fn_sym);
info->ancestor_entry_points.Add(ep_sym);
for (const auto& callee : caller_to_callee_[fn_sym]) { for (const auto& callee : caller_to_callee_[fn_sym]) {
set_entry_points(callee, ep_sym); set_entry_points(callee, ep_sym);
@ -166,9 +170,11 @@ bool TypeDeterminer::DetermineFunctions(const ast::FunctionList& funcs) {
} }
bool TypeDeterminer::DetermineFunction(ast::Function* func) { bool TypeDeterminer::DetermineFunction(ast::Function* func) {
symbol_to_function_[func->symbol()] = func; auto* info = function_infos_.Create<FunctionInfo>(func);
symbol_to_function_[func->symbol()] = info;
function_to_info_.emplace(func, info);
current_function_ = func; current_function_ = info;
variable_stack_.push_scope(); variable_stack_.push_scope();
for (auto* param : func->params()) { for (auto* param : func->params()) {
@ -409,19 +415,20 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
} }
} else { } else {
if (current_function_) { if (current_function_) {
caller_to_callee_[current_function_->symbol()].push_back( caller_to_callee_[current_function_->declaration->symbol()].push_back(
ident->symbol()); ident->symbol());
auto* callee_func = builder_->AST().Functions().Find(ident->symbol()); auto callee_func_it = symbol_to_function_.find(ident->symbol());
if (callee_func == nullptr) { if (callee_func_it == symbol_to_function_.end()) {
set_error(expr->source(), set_error(expr->source(),
"unable to find called function: " + "unable to find called function: " +
builder_->Symbols().NameFor(ident->symbol())); builder_->Symbols().NameFor(ident->symbol()));
return false; return false;
} }
auto* callee_func = callee_func_it->second;
// We inherit any referenced variables from the callee. // We inherit any referenced variables from the callee.
for (auto* var : callee_func->referenced_module_variables()) { for (auto* var : callee_func->referenced_module_vars) {
set_referenced_from_function_if_needed(var, false); set_referenced_from_function_if_needed(var, false);
} }
} }
@ -828,7 +835,7 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
auto iter = symbol_to_function_.find(symbol); auto iter = symbol_to_function_.find(symbol);
if (iter != symbol_to_function_.end()) { if (iter != symbol_to_function_.end()) {
SetType(expr, iter->second->return_type()); SetType(expr, iter->second->declaration->return_type());
return true; return true;
} }
@ -1204,4 +1211,25 @@ void TypeDeterminer::SetType(ast::Expression* expr, type::Type* type) const {
builder_->create<semantic::Expression>(type)); builder_->create<semantic::Expression>(type));
} }
void TypeDeterminer::CreateSemanticFunctions() const {
for (auto it : function_to_info_) {
auto* func = it.first;
auto* info = it.second;
if (builder_->Sem().Get(func)) {
// ast::Function already has a semantic::Function node.
// This is likely via explicit call to DetermineXXX() in test.
continue;
}
builder_->Sem().Add(func, builder_->create<semantic::Function>(
info->referenced_module_vars,
info->local_referenced_module_vars,
info->ancestor_entry_points));
}
}
TypeDeterminer::FunctionInfo::FunctionInfo(ast::Function* decl)
: declaration(decl) {}
TypeDeterminer::FunctionInfo::~FunctionInfo() = default;
} // namespace tint } // namespace tint

42
src/type_determiner.h
View File

@ -17,6 +17,7 @@
#include <string> #include <string>
#include <unordered_map> #include <unordered_map>
#include <unordered_set>
#include <vector> #include <vector>
#include "src/ast/module.h" #include "src/ast/module.h"
@ -98,6 +99,10 @@ class TypeDeterminer {
/// @returns false on error /// @returns false on error
bool DetermineStorageTextureSubtype(type::StorageTexture* tex); bool DetermineStorageTextureSubtype(type::StorageTexture* tex);
/// Creates the semantic::Function nodes and adds them to the semantic::Info
/// of the ProgramBuilder.
void CreateSemanticFunctions() const;
/// Testing method to set a given variable into the type stack /// Testing method to set a given variable into the type stack
/// @param var the variable to set /// @param var the variable to set
void RegisterVariableForTesting(ast::Variable* var) { void RegisterVariableForTesting(ast::Variable* var) {
@ -123,6 +128,37 @@ class TypeDeterminer {
bool SetIntrinsicIfNeeded(ast::IdentifierExpression* ident); bool SetIntrinsicIfNeeded(ast::IdentifierExpression* ident);
private: private:
template <typename T>
struct UniqueVector {
using ConstIterator = typename std::vector<T>::const_iterator;
void Add(const T& val) {
if (set.count(val) == 0) {
vector.emplace_back(val);
set.emplace(val);
}
}
ConstIterator begin() const { return vector.begin(); }
ConstIterator end() const { return vector.end(); }
operator const std::vector<T> &() const { return vector; }
private:
std::vector<T> vector;
std::unordered_set<T> set;
};
/// Structure holding semantic information about a function.
/// Used to build the semantic::Function nodes at the end of resolving.
struct FunctionInfo {
explicit FunctionInfo(ast::Function* decl);
~FunctionInfo();
ast::Function* const declaration;
UniqueVector<ast::Variable*> referenced_module_vars;
UniqueVector<ast::Variable*> local_referenced_module_vars;
UniqueVector<Symbol> ancestor_entry_points;
};
void set_error(const Source& src, const std::string& msg); void set_error(const Source& src, const std::string& msg);
void set_referenced_from_function_if_needed(ast::Variable* var, bool local); void set_referenced_from_function_if_needed(ast::Variable* var, bool local);
void set_entry_points(const Symbol& fn_sym, Symbol ep_sym); void set_entry_points(const Symbol& fn_sym, Symbol ep_sym);
@ -153,8 +189,10 @@ class TypeDeterminer {
ProgramBuilder* builder_; ProgramBuilder* builder_;
std::string error_; std::string error_;
ScopeStack<ast::Variable*> variable_stack_; ScopeStack<ast::Variable*> variable_stack_;
std::unordered_map<Symbol, ast::Function*> symbol_to_function_; std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_;
ast::Function* current_function_ = nullptr; std::unordered_map<ast::Function*, FunctionInfo*> function_to_info_;
FunctionInfo* current_function_ = nullptr;
BlockAllocator<FunctionInfo> function_infos_;
// Map from caller functions to callee functions. // Map from caller functions to callee functions.
std::unordered_map<Symbol, std::vector<Symbol>> caller_to_callee_; std::unordered_map<Symbol, std::vector<Symbol>> caller_to_callee_;

37
src/type_determiner_test.cc
View File

@ -52,6 +52,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/program_builder.h" #include "src/program_builder.h"
#include "src/semantic/expression.h" #include "src/semantic/expression.h"
#include "src/semantic/function.h"
#include "src/type/alias_type.h" #include "src/type/alias_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
#include "src/type/bool_type.h" #include "src/type/bool_type.h"
@ -659,7 +660,10 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables) {
// Register the function // Register the function
EXPECT_TRUE(td()->Determine()); EXPECT_TRUE(td()->Determine());
const auto& vars = func->referenced_module_variables(); auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
const auto& vars = func_sem->ReferencedModuleVariables();
ASSERT_EQ(vars.size(), 5u); ASSERT_EQ(vars.size(), 5u);
EXPECT_EQ(vars[0], out_var); EXPECT_EQ(vars[0], out_var);
EXPECT_EQ(vars[1], in_var); EXPECT_EQ(vars[1], in_var);
@ -700,7 +704,10 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) {
// Register the function // Register the function
EXPECT_TRUE(td()->Determine()); EXPECT_TRUE(td()->Determine());
const auto& vars = func2->referenced_module_variables(); auto* func2_sem = Sem().Get(func2);
ASSERT_NE(func2_sem, nullptr);
const auto& vars = func2_sem->ReferencedModuleVariables();
ASSERT_EQ(vars.size(), 5u); ASSERT_EQ(vars.size(), 5u);
EXPECT_EQ(vars[0], out_var); EXPECT_EQ(vars[0], out_var);
EXPECT_EQ(vars[1], in_var); EXPECT_EQ(vars[1], in_var);
@ -726,7 +733,10 @@ TEST_F(TypeDeterminerTest, Function_NotRegisterFunctionVariable) {
// Register the function // Register the function
EXPECT_TRUE(td()->Determine()) << td()->error(); EXPECT_TRUE(td()->Determine()) << td()->error();
EXPECT_EQ(func->referenced_module_variables().size(), 0u); auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->ReferencedModuleVariables().size(), 0u);
} }
TEST_F(TypeDeterminerTest, Expr_MemberAccessor_Struct) { TEST_F(TypeDeterminerTest, Expr_MemberAccessor_Struct) {
@ -2284,22 +2294,33 @@ TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) {
// Register the functions and calculate the callers // Register the functions and calculate the callers
ASSERT_TRUE(td()->Determine()) << td()->error(); ASSERT_TRUE(td()->Determine()) << td()->error();
const auto& b_eps = func_b->ancestor_entry_points(); auto* func_b_sem = Sem().Get(func_b);
auto* func_a_sem = Sem().Get(func_a);
auto* func_c_sem = Sem().Get(func_c);
auto* ep_1_sem = Sem().Get(ep_1);
auto* ep_2_sem = Sem().Get(ep_2);
ASSERT_NE(func_b_sem, nullptr);
ASSERT_NE(func_a_sem, nullptr);
ASSERT_NE(func_c_sem, nullptr);
ASSERT_NE(ep_1_sem, nullptr);
ASSERT_NE(ep_2_sem, nullptr);
const auto& b_eps = func_b_sem->AncestorEntryPoints();
ASSERT_EQ(2u, b_eps.size()); ASSERT_EQ(2u, b_eps.size());
EXPECT_EQ(Symbols().Register("ep_1"), b_eps[0]); EXPECT_EQ(Symbols().Register("ep_1"), b_eps[0]);
EXPECT_EQ(Symbols().Register("ep_2"), b_eps[1]); EXPECT_EQ(Symbols().Register("ep_2"), b_eps[1]);
const auto& a_eps = func_a->ancestor_entry_points(); const auto& a_eps = func_a_sem->AncestorEntryPoints();
ASSERT_EQ(1u, a_eps.size()); ASSERT_EQ(1u, a_eps.size());
EXPECT_EQ(Symbols().Register("ep_1"), a_eps[0]); EXPECT_EQ(Symbols().Register("ep_1"), a_eps[0]);
const auto& c_eps = func_c->ancestor_entry_points(); const auto& c_eps = func_c_sem->AncestorEntryPoints();
ASSERT_EQ(2u, c_eps.size()); ASSERT_EQ(2u, c_eps.size());
EXPECT_EQ(Symbols().Register("ep_1"), c_eps[0]); EXPECT_EQ(Symbols().Register("ep_1"), c_eps[0]);
EXPECT_EQ(Symbols().Register("ep_2"), c_eps[1]); EXPECT_EQ(Symbols().Register("ep_2"), c_eps[1]);
EXPECT_TRUE(ep_1->ancestor_entry_points().empty()); EXPECT_TRUE(ep_1_sem->AncestorEntryPoints().empty());
EXPECT_TRUE(ep_2->ancestor_entry_points().empty()); EXPECT_TRUE(ep_2_sem->AncestorEntryPoints().empty());
} }
using TypeDeterminerTextureIntrinsicTest = using TypeDeterminerTextureIntrinsicTest =

60
src/writer/hlsl/generator_impl.cc
View File

@ -46,6 +46,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/program_builder.h" #include "src/program_builder.h"
#include "src/semantic/expression.h" #include "src/semantic/expression.h"
#include "src/semantic/function.h"
#include "src/type/access_control_type.h" #include "src/type/access_control_type.h"
#include "src/type/alias_type.h" #include "src/type/alias_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
@ -591,15 +592,17 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
out << name << "("; out << name << "(";
auto* func_sem = builder_.Sem().Get(func);
bool first = true; bool first = true;
if (has_referenced_in_var_needing_struct(func)) { if (has_referenced_in_var_needing_struct(func_sem)) {
auto var_name = current_ep_var_name(VarType::kIn); auto var_name = current_ep_var_name(VarType::kIn);
if (!var_name.empty()) { if (!var_name.empty()) {
out << var_name; out << var_name;
first = false; first = false;
} }
} }
if (has_referenced_out_var_needing_struct(func)) { if (has_referenced_out_var_needing_struct(func_sem)) {
auto var_name = current_ep_var_name(VarType::kOut); auto var_name = current_ep_var_name(VarType::kOut);
if (!var_name.empty()) { if (!var_name.empty()) {
if (!first) { if (!first) {
@ -1223,15 +1226,16 @@ bool GeneratorImpl::EmitIf(std::ostream& out, ast::IfStatement* stmt) {
return true; return true;
} }
bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) { bool GeneratorImpl::has_referenced_in_var_needing_struct(
for (auto data : func->referenced_location_variables()) { const semantic::Function* func) {
for (auto data : func->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kInput) { if (var->storage_class() == ast::StorageClass::kInput) {
return true; return true;
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kInput) { if (var->storage_class() == ast::StorageClass::kInput) {
return true; return true;
@ -1240,15 +1244,16 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) {
return false; return false;
} }
bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) { bool GeneratorImpl::has_referenced_out_var_needing_struct(
for (auto data : func->referenced_location_variables()) { const semantic::Function* func) {
for (auto data : func->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput) { if (var->storage_class() == ast::StorageClass::kOutput) {
return true; return true;
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput) { if (var->storage_class() == ast::StorageClass::kOutput) {
return true; return true;
@ -1257,8 +1262,9 @@ bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) {
return false; return false;
} }
bool GeneratorImpl::has_referenced_var_needing_struct(ast::Function* func) { bool GeneratorImpl::has_referenced_var_needing_struct(
for (auto data : func->referenced_location_variables()) { const semantic::Function* func) {
for (auto data : func->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput || if (var->storage_class() == ast::StorageClass::kOutput ||
var->storage_class() == ast::StorageClass::kInput) { var->storage_class() == ast::StorageClass::kInput) {
@ -1266,7 +1272,7 @@ bool GeneratorImpl::has_referenced_var_needing_struct(ast::Function* func) {
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput || if (var->storage_class() == ast::StorageClass::kOutput ||
var->storage_class() == ast::StorageClass::kInput) { var->storage_class() == ast::StorageClass::kInput) {
@ -1284,14 +1290,16 @@ bool GeneratorImpl::EmitFunction(std::ostream& out, ast::Function* func) {
return true; return true;
} }
auto* func_sem = builder_.Sem().Get(func);
// TODO(dsinclair): This could be smarter. If the input/outputs for multiple // TODO(dsinclair): This could be smarter. If the input/outputs for multiple
// entry points are the same we could generate a single struct and then have // entry points are the same we could generate a single struct and then have
// this determine it's the same struct and just emit once. // this determine it's the same struct and just emit once.
bool emit_duplicate_functions = func->ancestor_entry_points().size() > 0 && bool emit_duplicate_functions = func_sem->AncestorEntryPoints().size() > 0 &&
has_referenced_var_needing_struct(func); has_referenced_var_needing_struct(func_sem);
if (emit_duplicate_functions) { if (emit_duplicate_functions) {
for (const auto& ep_sym : func->ancestor_entry_points()) { for (const auto& ep_sym : func_sem->AncestorEntryPoints()) {
if (!EmitFunctionInternal(out, func, emit_duplicate_functions, ep_sym)) { if (!EmitFunctionInternal(out, func, emit_duplicate_functions, ep_sym)) {
return false; return false;
} }
@ -1394,7 +1402,10 @@ bool GeneratorImpl::EmitEntryPointData(
std::unordered_set<Symbol>& emitted_globals) { std::unordered_set<Symbol>& emitted_globals) {
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> in_variables; std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> in_variables;
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> outvariables; std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> outvariables;
for (auto data : func->referenced_location_variables()) { auto* func_sem = builder_.Sem().Get(func);
auto func_sym = func->symbol();
for (auto data : func_sem->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
auto* deco = data.second; auto* deco = data.second;
@ -1405,7 +1416,7 @@ bool GeneratorImpl::EmitEntryPointData(
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
auto* deco = data.second; auto* deco = data.second;
@ -1417,7 +1428,7 @@ bool GeneratorImpl::EmitEntryPointData(
} }
bool emitted_uniform = false; bool emitted_uniform = false;
for (auto data : func->referenced_uniform_variables()) { for (auto data : func_sem->ReferencedUniformVariables()) {
auto* var = data.first; auto* var = data.first;
// TODO(dsinclair): We're using the binding to make up the buffer number but // TODO(dsinclair): We're using the binding to make up the buffer number but
// we should instead be using a provided mapping that uses both buffer and // we should instead be using a provided mapping that uses both buffer and
@ -1471,7 +1482,7 @@ bool GeneratorImpl::EmitEntryPointData(
} }
bool emitted_storagebuffer = false; bool emitted_storagebuffer = false;
for (auto data : func->referenced_storagebuffer_variables()) { for (auto data : func_sem->ReferencedStoragebufferVariables()) {
auto* var = data.first; auto* var = data.first;
auto* binding = data.second.binding; auto* binding = data.second.binding;
@ -1500,10 +1511,10 @@ bool GeneratorImpl::EmitEntryPointData(
} }
if (!in_variables.empty()) { if (!in_variables.empty()) {
auto in_struct_name = generate_name( auto in_struct_name = generate_name(builder_.Symbols().NameFor(func_sym) +
builder_.Symbols().NameFor(func->symbol()) + "_" + kInStructNameSuffix); "_" + kInStructNameSuffix);
auto in_var_name = generate_name(kTintStructInVarPrefix); auto in_var_name = generate_name(kTintStructInVarPrefix);
ep_sym_to_in_data_[func->symbol()] = {in_struct_name, in_var_name}; ep_sym_to_in_data_[func_sym] = {in_struct_name, in_var_name};
make_indent(out); make_indent(out);
out << "struct " << in_struct_name << " {" << std::endl; out << "struct " << in_struct_name << " {" << std::endl;
@ -1547,11 +1558,10 @@ bool GeneratorImpl::EmitEntryPointData(
} }
if (!outvariables.empty()) { if (!outvariables.empty()) {
auto outstruct_name = auto outstruct_name = generate_name(builder_.Symbols().NameFor(func_sym) +
generate_name(builder_.Symbols().NameFor(func->symbol()) + "_" + "_" + kOutStructNameSuffix);
kOutStructNameSuffix);
auto outvar_name = generate_name(kTintStructOutVarPrefix); auto outvar_name = generate_name(kTintStructOutVarPrefix);
ep_sym_to_out_data_[func->symbol()] = {outstruct_name, outvar_name}; ep_sym_to_out_data_[func_sym] = {outstruct_name, outvar_name};
make_indent(out); make_indent(out);
out << "struct " << outstruct_name << " {" << std::endl; out << "struct " << outstruct_name << " {" << std::endl;

6
src/writer/hlsl/generator_impl.h
View File

@ -356,16 +356,16 @@ class GeneratorImpl {
/// Determines if the function needs the input struct passed to it. /// Determines if the function needs the input struct passed to it.
/// @param func the function to check /// @param func the function to check
/// @returns true if there are input struct variables used in the function /// @returns true if there are input struct variables used in the function
bool has_referenced_in_var_needing_struct(ast::Function* func); bool has_referenced_in_var_needing_struct(const semantic::Function* func);
/// Determines if the function needs the output struct passed to it. /// Determines if the function needs the output struct passed to it.
/// @param func the function to check /// @param func the function to check
/// @returns true if there are output struct variables used in the function /// @returns true if there are output struct variables used in the function
bool has_referenced_out_var_needing_struct(ast::Function* func); bool has_referenced_out_var_needing_struct(const semantic::Function* func);
/// Determines if any used program variable requires an input or output /// Determines if any used program variable requires an input or output
/// struct. /// struct.
/// @param func the function to check /// @param func the function to check
/// @returns true if an input or output struct is required. /// @returns true if an input or output struct is required.
bool has_referenced_var_needing_struct(ast::Function* func); bool has_referenced_var_needing_struct(const semantic::Function* func);
/// @returns the namer for testing /// @returns the namer for testing
Namer* namer_for_testing() { return &namer_; } Namer* namer_for_testing() { return &namer_; }

46
src/writer/msl/generator_impl.cc
View File

@ -51,6 +51,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/program.h" #include "src/program.h"
#include "src/semantic/expression.h" #include "src/semantic/expression.h"
#include "src/semantic/function.h"
#include "src/type/access_control_type.h" #include "src/type/access_control_type.h"
#include "src/type/alias_type.h" #include "src/type/alias_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
@ -504,7 +505,8 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
} }
} }
for (const auto& data : func->referenced_builtin_variables()) { auto* func_sem = program_->Sem().Get(func);
for (const auto& data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) { if (var->storage_class() != ast::StorageClass::kInput) {
continue; continue;
@ -516,7 +518,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
out_ << program_->Symbols().NameFor(var->symbol()); out_ << program_->Symbols().NameFor(var->symbol());
} }
for (const auto& data : func->referenced_uniform_variables()) { for (const auto& data : func_sem->ReferencedUniformVariables()) {
auto* var = data.first; auto* var = data.first;
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
@ -525,7 +527,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
out_ << program_->Symbols().NameFor(var->symbol()); out_ << program_->Symbols().NameFor(var->symbol());
} }
for (const auto& data : func->referenced_storagebuffer_variables()) { for (const auto& data : func_sem->ReferencedStoragebufferVariables()) {
auto* var = data.first; auto* var = data.first;
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
@ -1021,10 +1023,13 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
} }
bool GeneratorImpl::EmitEntryPointData(ast::Function* func) { bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
auto* func_sem = program_->Sem().Get(func);
std::vector<std::pair<ast::Variable*, uint32_t>> in_locations; std::vector<std::pair<ast::Variable*, uint32_t>> in_locations;
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>>
out_variables; out_variables;
for (auto data : func->referenced_location_variables()) {
for (auto data : func_sem->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
auto* deco = data.second; auto* deco = data.second;
@ -1035,7 +1040,7 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
auto* deco = data.second; auto* deco = data.second;
@ -1183,7 +1188,8 @@ void GeneratorImpl::EmitStage(ast::PipelineStage stage) {
} }
bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) { bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) {
for (auto data : func->referenced_location_variables()) { auto* func_sem = program_->Sem().Get(func);
for (auto data : func_sem->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kInput) { if (var->storage_class() == ast::StorageClass::kInput) {
return true; return true;
@ -1193,14 +1199,16 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) {
} }
bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) { bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) {
for (auto data : func->referenced_location_variables()) { auto* func_sem = program_->Sem().Get(func);
for (auto data : func_sem->ReferencedLocationVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput) { if (var->storage_class() == ast::StorageClass::kOutput) {
return true; return true;
} }
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() == ast::StorageClass::kOutput) { if (var->storage_class() == ast::StorageClass::kOutput) {
return true; return true;
@ -1215,6 +1223,8 @@ bool GeneratorImpl::has_referenced_var_needing_struct(ast::Function* func) {
} }
bool GeneratorImpl::EmitFunction(ast::Function* func) { bool GeneratorImpl::EmitFunction(ast::Function* func) {
auto* func_sem = program_->Sem().Get(func);
make_indent(); make_indent();
// Entry points will be emitted later, skip for now. // Entry points will be emitted later, skip for now.
@ -1225,11 +1235,11 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
// TODO(dsinclair): This could be smarter. If the input/outputs for multiple // TODO(dsinclair): This could be smarter. If the input/outputs for multiple
// entry points are the same we could generate a single struct and then have // entry points are the same we could generate a single struct and then have
// this determine it's the same struct and just emit once. // this determine it's the same struct and just emit once.
bool emit_duplicate_functions = func->ancestor_entry_points().size() > 0 && bool emit_duplicate_functions = func_sem->AncestorEntryPoints().size() > 0 &&
has_referenced_var_needing_struct(func); has_referenced_var_needing_struct(func);
if (emit_duplicate_functions) { if (emit_duplicate_functions) {
for (const auto& ep_sym : func->ancestor_entry_points()) { for (const auto& ep_sym : func_sem->AncestorEntryPoints()) {
if (!EmitFunctionInternal(func, emit_duplicate_functions, ep_sym)) { if (!EmitFunctionInternal(func, emit_duplicate_functions, ep_sym)) {
return false; return false;
} }
@ -1249,6 +1259,8 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
bool emit_duplicate_functions, bool emit_duplicate_functions,
Symbol ep_sym) { Symbol ep_sym) {
auto* func_sem = program_->Sem().Get(func);
auto name = func->symbol().to_str(); auto name = func->symbol().to_str();
if (!EmitType(func->return_type(), "")) { if (!EmitType(func->return_type(), "")) {
return false; return false;
@ -1294,7 +1306,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
} }
} }
for (const auto& data : func->referenced_builtin_variables()) { for (const auto& data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) { if (var->storage_class() != ast::StorageClass::kInput) {
continue; continue;
@ -1311,7 +1323,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
out_ << "& " << program_->Symbols().NameFor(var->symbol()); out_ << "& " << program_->Symbols().NameFor(var->symbol());
} }
for (const auto& data : func->referenced_uniform_variables()) { for (const auto& data : func_sem->ReferencedUniformVariables()) {
auto* var = data.first; auto* var = data.first;
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
@ -1326,7 +1338,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
out_ << "& " << program_->Symbols().NameFor(var->symbol()); out_ << "& " << program_->Symbols().NameFor(var->symbol());
} }
for (const auto& data : func->referenced_storagebuffer_variables()) { for (const auto& data : func_sem->ReferencedStoragebufferVariables()) {
auto* var = data.first; auto* var = data.first;
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
@ -1404,6 +1416,8 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const {
} }
bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
auto* func_sem = program_->Sem().Get(func);
make_indent(); make_indent();
current_ep_sym_ = func->symbol(); current_ep_sym_ = func->symbol();
@ -1431,7 +1445,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
first = false; first = false;
} }
for (auto data : func->referenced_builtin_variables()) { for (auto data : func_sem->ReferencedBuiltinVariables()) {
auto* var = data.first; auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) { if (var->storage_class() != ast::StorageClass::kInput) {
continue; continue;
@ -1457,7 +1471,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
<< "]]"; << "]]";
} }
for (auto data : func->referenced_uniform_variables()) { for (auto data : func_sem->ReferencedUniformVariables()) {
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
} }
@ -1485,7 +1499,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
<< binding->value() << ")]]"; << binding->value() << ")]]";
} }
for (auto data : func->referenced_storagebuffer_variables()) { for (auto data : func_sem->ReferencedStoragebufferVariables()) {
if (!first) { if (!first) {
out_ << ", "; out_ << ", ";
} }

7
src/writer/spirv/builder.cc
View File

@ -60,6 +60,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/program.h" #include "src/program.h"
#include "src/semantic/expression.h" #include "src/semantic/expression.h"
#include "src/semantic/function.h"
#include "src/type/access_control_type.h" #include "src/type/access_control_type.h"
#include "src/type/alias_type.h" #include "src/type/alias_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
@ -457,7 +458,8 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
Operand::Int(stage), Operand::Int(id), Operand::Int(stage), Operand::Int(id),
Operand::String(builder_.Symbols().NameFor(func->symbol()))}; Operand::String(builder_.Symbols().NameFor(func->symbol()))};
for (const auto* var : func->referenced_module_variables()) { auto* func_sem = builder_.Sem().Get(func);
for (const auto* var : func_sem->ReferencedModuleVariables()) {
// For SPIR-V 1.3 we only output Input/output variables. If we update to // For SPIR-V 1.3 we only output Input/output variables. If we update to
// SPIR-V 1.4 or later this should be all variables. // SPIR-V 1.4 or later this should be all variables.
if (var->storage_class() != ast::StorageClass::kInput && if (var->storage_class() != ast::StorageClass::kInput &&
@ -496,7 +498,8 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
Operand::Int(x), Operand::Int(y), Operand::Int(z)}); Operand::Int(x), Operand::Int(y), Operand::Int(z)});
} }
for (auto builtin : func->referenced_builtin_variables()) { auto* func_sem = builder_.Sem().Get(func);
for (auto builtin : func_sem->ReferencedBuiltinVariables()) {
if (builtin.second->value() == ast::Builtin::kFragDepth) { if (builtin.second->value() == ast::Builtin::kFragDepth) {
push_execution_mode( push_execution_mode(
spv::Op::OpExecutionMode, spv::Op::OpExecutionMode,

9
src/writer/spirv/builder_function_decoration_test.cc
View File

@ -24,6 +24,7 @@
#include "src/ast/stage_decoration.h" #include "src/ast/stage_decoration.h"
#include "src/ast/variable.h" #include "src/ast/variable.h"
#include "src/ast/workgroup_decoration.h" #include "src/ast/workgroup_decoration.h"
#include "src/semantic/function.h"
#include "src/type_determiner.h" #include "src/type_determiner.h"
#include "src/writer/spirv/builder.h" #include "src/writer/spirv/builder.h"
#include "src/writer/spirv/spv_dump.h" #include "src/writer/spirv/spv_dump.h"
@ -153,12 +154,6 @@ TEST_F(BuilderTest, FunctionDecoration_Stage_WithUsedInterfaceIds) {
AST().AddGlobalVariable(v_out); AST().AddGlobalVariable(v_out);
AST().AddGlobalVariable(v_wg); AST().AddGlobalVariable(v_wg);
td.RegisterVariableForTesting(v_in);
td.RegisterVariableForTesting(v_out);
td.RegisterVariableForTesting(v_wg);
ASSERT_TRUE(td.DetermineFunction(func)) << td.error();
spirv::Builder& b = Build(); spirv::Builder& b = Build();
EXPECT_TRUE(b.GenerateGlobalVariable(v_in)) << b.error(); EXPECT_TRUE(b.GenerateGlobalVariable(v_in)) << b.error();
@ -285,8 +280,6 @@ TEST_F(BuilderTest, FunctionDecoration_ExecutionMode_FragDepth) {
}, },
ast::FunctionDecorationList{}); ast::FunctionDecorationList{});
func->add_referenced_module_variable(fragdepth);
spirv::Builder& b = Build(); spirv::Builder& b = Build();
ASSERT_TRUE(b.GenerateExecutionModes(func, 3)) << b.error(); ASSERT_TRUE(b.GenerateExecutionModes(func, 3)) << b.error();

6
src/writer/wgsl/generator_impl.cc
View File

@ -58,6 +58,7 @@
#include "src/ast/variable_decl_statement.h" #include "src/ast/variable_decl_statement.h"
#include "src/ast/workgroup_decoration.h" #include "src/ast/workgroup_decoration.h"
#include "src/program.h" #include "src/program.h"
#include "src/semantic/function.h"
#include "src/type/access_control_type.h" #include "src/type/access_control_type.h"
#include "src/type/alias_type.h" #include "src/type/alias_type.h"
#include "src/type/array_type.h" #include "src/type/array_type.h"
@ -146,7 +147,7 @@ bool GeneratorImpl::GenerateEntryPoint(ast::PipelineStage stage,
} }
bool found_func_variable = false; bool found_func_variable = false;
for (auto* var : func->referenced_module_variables()) { for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) {
if (!EmitVariable(var)) { if (!EmitVariable(var)) {
return false; return false;
} }
@ -157,7 +158,8 @@ bool GeneratorImpl::GenerateEntryPoint(ast::PipelineStage stage,
} }
for (auto* f : program_->AST().Functions()) { for (auto* f : program_->AST().Functions()) {
if (!f->HasAncestorEntryPoint(program_->Symbols().Get(name))) { auto* f_sem = program_->Sem().Get(f);
if (!f_sem->HasAncestorEntryPoint(program_->Symbols().Get(name))) {
continue; continue;
} }