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

sem: Rename Function methods 

Explicitly name references as either Direct or Transitive.
Rename workgroup_size() to WorkgroupSize().
Remove the return_statements. These were not used anywhere.

Change-Id: I7191665db9c3211d086dd90939abec7003cd7be7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68405
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton 2021-11-04 22:29:22 +00:00
parent cc26a4166d
commit 2423df3e04
12 changed files with 169 additions and 189 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

30
src/inspector/inspector.cc
View File

@ -145,7 +145,7 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
entry_point.remapped_name = program_->Symbols().NameFor(func->symbol);
entry_point.stage = func->PipelineStage();
auto wgsize = sem->workgroup_size();
auto wgsize = sem->WorkgroupSize();
entry_point.workgroup_size_x = wgsize[0].value;
entry_point.workgroup_size_y = wgsize[1].value;
entry_point.workgroup_size_z = wgsize[2].value;
@ -188,7 +188,7 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
func->return_type_decorations);
}
for (auto* var : sem->ReferencedModuleVariables()) {
for (auto* var : sem->TransitivelyReferencedGlobals()) {
auto* decl = var->Declaration();
auto name = program_->Symbols().NameFor(decl->symbol);
@ -333,13 +333,13 @@ uint32_t Inspector::GetStorageSize(const std::string& entry_point) {
size_t size = 0;
auto* func_sem = program_->Sem().Get(func);
for (auto& ruv : func_sem->ReferencedUniformVariables()) {
for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
const sem::Struct* s = ruv.first->Type()->UnwrapRef()->As<sem::Struct>();
if (s && s->IsBlockDecorated()) {
size += s->Size();
}
}
for (auto& rsv : func_sem->ReferencedStorageBufferVariables()) {
for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
const sem::Struct* s = rsv.first->Type()->UnwrapRef()->As<sem::Struct>();
if (s) {
size += s->Size();
@ -389,7 +389,7 @@ std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
for (auto& ruv : func_sem->ReferencedUniformVariables()) {
for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
auto* var = ruv.first;
auto binding_info = ruv.second;
@ -437,7 +437,7 @@ std::vector<ResourceBinding> Inspector::GetSamplerResourceBindings(
std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
for (auto& rs : func_sem->ReferencedSamplerVariables()) {
for (auto& rs : func_sem->TransitivelyReferencedSamplerVariables()) {
auto binding_info = rs.second;
ResourceBinding entry;
@ -461,7 +461,8 @@ std::vector<ResourceBinding> Inspector::GetComparisonSamplerResourceBindings(
std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
for (auto& rcs : func_sem->ReferencedComparisonSamplerVariables()) {
for (auto& rcs :
func_sem->TransitivelyReferencedComparisonSamplerVariables()) {
auto binding_info = rcs.second;
ResourceBinding entry;
@ -502,7 +503,8 @@ std::vector<ResourceBinding> Inspector::GetTextureResourceBindings(
std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
for (auto& ref : func_sem->ReferencedVariablesOfType(texture_type)) {
for (auto& ref :
func_sem->TransitivelyReferencedVariablesOfType(texture_type)) {
auto* var = ref.first;
auto binding_info = ref.second;
@ -567,7 +569,7 @@ uint32_t Inspector::GetWorkgroupStorageSize(const std::string& entry_point) {
uint32_t total_size = 0;
auto* func_sem = program_->Sem().Get(func);
for (const sem::Variable* var : func_sem->ReferencedModuleVariables()) {
for (const sem::Variable* var : func_sem->TransitivelyReferencedGlobals()) {
if (var->StorageClass() == ast::StorageClass::kWorkgroup) {
auto* ty = var->Type()->UnwrapRef();
uint32_t align = ty->Align();
@ -678,7 +680,7 @@ std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
auto* func_sem = program_->Sem().Get(func);
std::vector<ResourceBinding> result;
for (auto& rsv : func_sem->ReferencedStorageBufferVariables()) {
for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
auto* var = rsv.first;
auto binding_info = rsv.second;
@ -717,8 +719,9 @@ std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
std::vector<ResourceBinding> result;
auto* func_sem = program_->Sem().Get(func);
auto referenced_variables =
multisampled_only ? func_sem->ReferencedMultisampledTextureVariables()
: func_sem->ReferencedSampledTextureVariables();
multisampled_only
? func_sem->TransitivelyReferencedMultisampledTextureVariables()
: func_sem->TransitivelyReferencedSampledTextureVariables();
for (auto& ref : referenced_variables) {
auto* var = ref.first;
auto binding_info = ref.second;
@ -757,7 +760,8 @@ std::vector<ResourceBinding> Inspector::GetStorageTextureResourceBindingsImpl(
auto* func_sem = program_->Sem().Get(func);
std::vector<ResourceBinding> result;
for (auto& ref : func_sem->ReferencedVariablesOfType<sem::StorageTexture>()) {
for (auto& ref :
func_sem->TransitivelyReferencedVariablesOfType<sem::StorageTexture>()) {
auto* var = ref.first;
auto binding_info = ref.second;

9
src/resolver/resolver.cc
View File

@ -3795,10 +3795,10 @@ void Resolver::CreateSemanticNodes() const {
}
auto remap_vars = [&sem](const std::vector<VariableInfo*>& in) {
std::vector<const sem::Variable*> out;
std::vector<const sem::GlobalVariable*> out;
out.reserve(in.size());
for (auto* info : in) {
out.emplace_back(sem.Get(info->declaration));
out.emplace_back(sem.Get<sem::GlobalVariable>(info->declaration));
}
return out;
};
@ -3818,9 +3818,8 @@ void Resolver::CreateSemanticNodes() const {
auto* sem_func = builder_->create<sem::Function>(
info->declaration, info->return_type, parameters,
remap_vars(info->referenced_module_vars),
remap_vars(info->local_referenced_module_vars), info->return_statements,
info->callsites, ancestor_entry_points[func->symbol],
info->workgroup_size);
remap_vars(info->local_referenced_module_vars), info->callsites,
ancestor_entry_points[func->symbol], info->workgroup_size);
func_info_to_sem_func.emplace(info, sem_func);
sem.Add(func, sem_func);
}

118
src/resolver/resolver_test.cc
View File

@ -819,7 +819,7 @@ TEST_F(ResolverTest, Function_RegisterInputOutputVariables) {
EXPECT_EQ(func_sem->Parameters().size(), 0u);
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::Void>());
const auto& vars = func_sem->ReferencedModuleVariables();
const auto& vars = func_sem->TransitivelyReferencedGlobals();
ASSERT_EQ(vars.size(), 3u);
EXPECT_EQ(vars[0]->Declaration(), wg_var);
EXPECT_EQ(vars[1]->Declaration(), sb_var);
@ -856,7 +856,7 @@ TEST_F(ResolverTest, Function_RegisterInputOutputVariables_SubFunction) {
ASSERT_NE(func2_sem, nullptr);
EXPECT_EQ(func2_sem->Parameters().size(), 0u);
const auto& vars = func2_sem->ReferencedModuleVariables();
const auto& vars = func2_sem->TransitivelyReferencedGlobals();
ASSERT_EQ(vars.size(), 3u);
EXPECT_EQ(vars[0]->Declaration(), wg_var);
EXPECT_EQ(vars[1]->Declaration(), sb_var);
@ -875,7 +875,7 @@ TEST_F(ResolverTest, Function_NotRegisterFunctionVariable) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->ReferencedModuleVariables().size(), 0u);
EXPECT_EQ(func_sem->TransitivelyReferencedGlobals().size(), 0u);
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::Void>());
}
@ -890,7 +890,7 @@ TEST_F(ResolverTest, Function_NotRegisterFunctionConstant) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->ReferencedModuleVariables().size(), 0u);
EXPECT_EQ(func_sem->TransitivelyReferencedGlobals().size(), 0u);
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::Void>());
}
@ -902,34 +902,10 @@ TEST_F(ResolverTest, Function_NotRegisterFunctionParams) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->ReferencedModuleVariables().size(), 0u);
EXPECT_EQ(func_sem->TransitivelyReferencedGlobals().size(), 0u);
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::Void>());
}
TEST_F(ResolverTest, Function_ReturnStatements) {
auto* var = Var("foo", ty.f32());
auto* ret_1 = Return(1.f);
auto* ret_foo = Return("foo");
auto* func = Func("my_func", ast::VariableList{}, ty.f32(),
{
Decl(var),
If(true, Block(ret_1)),
ret_foo,
});
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->Parameters().size(), 0u);
EXPECT_EQ(func_sem->ReturnStatements().size(), 2u);
EXPECT_EQ(func_sem->ReturnStatements()[0], ret_1);
EXPECT_EQ(func_sem->ReturnStatements()[1], ret_foo);
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::F32>());
}
TEST_F(ResolverTest, Function_CallSites) {
auto* foo = Func("foo", ast::VariableList{}, ty.void_(), {});
@ -964,12 +940,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_NotSet) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 1u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 1u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 1u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 1u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 1u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 1u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Literals) {
@ -984,12 +960,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Literals) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 8u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 2u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 3u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 2u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 3u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Consts) {
@ -1010,12 +986,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Consts) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 16u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 8u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 2u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 16u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 8u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 2u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Consts_NestedInitializer) {
@ -1036,12 +1012,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Consts_NestedInitializer) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 8u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 4u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 1u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 4u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 1u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
}
TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts) {
@ -1062,12 +1038,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 16u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 8u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 2u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, width);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, height);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, depth);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 16u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 8u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 2u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, width);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, depth);
}
TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts_NoInit) {
@ -1088,12 +1064,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts_NoInit) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 0u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 0u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 0u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, width);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, height);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, depth);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 0u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 0u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 0u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, width);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, depth);
}
TEST_F(ResolverTest, Function_WorkgroupSize_Mixed) {
@ -1112,12 +1088,12 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Mixed) {
auto* func_sem = Sem().Get(func);
ASSERT_NE(func_sem, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[0].value, 8u);
EXPECT_EQ(func_sem->workgroup_size()[1].value, 2u);
EXPECT_EQ(func_sem->workgroup_size()[2].value, 3u);
EXPECT_EQ(func_sem->workgroup_size()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->workgroup_size()[1].overridable_const, height);
EXPECT_EQ(func_sem->workgroup_size()[2].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 2u);
EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 3u);
EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
}
TEST_F(ResolverTest, Expr_MemberAccessor_Struct) {

78
src/sem/function.cc
View File

@ -28,35 +28,35 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Function);
namespace tint {
namespace sem {
Function::Function(const ast::Function* declaration,
Function::Function(
const ast::Function* declaration,
Type* return_type,
std::vector<Parameter*> parameters,
std::vector<const Variable*> referenced_module_vars,
std::vector<const Variable*> local_referenced_module_vars,
std::vector<const ast::ReturnStatement*> return_statements,
std::vector<const GlobalVariable*> transitively_referenced_globals,
std::vector<const GlobalVariable*> directly_referenced_globals,
std::vector<const ast::CallExpression*> callsites,
std::vector<Symbol> ancestor_entry_points,
std::array<WorkgroupDimension, 3> workgroup_size)
sem::WorkgroupSize workgroup_size)
: Base(return_type, utils::ToConstPtrVec(parameters)),
declaration_(declaration),
referenced_module_vars_(std::move(referenced_module_vars)),
local_referenced_module_vars_(std::move(local_referenced_module_vars)),
return_statements_(std::move(return_statements)),
workgroup_size_(std::move(workgroup_size)),
directly_referenced_globals_(std::move(directly_referenced_globals)),
transitively_referenced_globals_(
std::move(transitively_referenced_globals)),
callsites_(callsites),
ancestor_entry_points_(std::move(ancestor_entry_points)),
workgroup_size_(std::move(workgroup_size)) {
ancestor_entry_points_(std::move(ancestor_entry_points)) {
for (auto* parameter : parameters) {
parameter->SetOwner(this);
}
}
} // namespace sem
Function::~Function() = default;
std::vector<std::pair<const Variable*, const ast::LocationDecoration*>>
Function::ReferencedLocationVariables() const {
Function::TransitivelyReferencedLocationVariables() const {
std::vector<std::pair<const Variable*, const ast::LocationDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
for (auto* deco : var->Declaration()->decorations) {
if (auto* location = deco->As<ast::LocationDecoration>()) {
ret.push_back({var, location});
@ -67,10 +67,11 @@ Function::ReferencedLocationVariables() const {
return ret;
}
Function::VariableBindings Function::ReferencedUniformVariables() const {
Function::VariableBindings Function::TransitivelyReferencedUniformVariables()
const {
VariableBindings ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
if (var->StorageClass() != ast::StorageClass::kUniform) {
continue;
}
@ -82,10 +83,11 @@ Function::VariableBindings Function::ReferencedUniformVariables() const {
return ret;
}
Function::VariableBindings Function::ReferencedStorageBufferVariables() const {
Function::VariableBindings
Function::TransitivelyReferencedStorageBufferVariables() const {
VariableBindings ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
if (var->StorageClass() != ast::StorageClass::kStorage) {
continue;
}
@ -98,10 +100,10 @@ Function::VariableBindings Function::ReferencedStorageBufferVariables() const {
}
std::vector<std::pair<const Variable*, const ast::BuiltinDecoration*>>
Function::ReferencedBuiltinVariables() const {
Function::TransitivelyReferencedBuiltinVariables() const {
std::vector<std::pair<const Variable*, const ast::BuiltinDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
for (auto* deco : var->Declaration()->decorations) {
if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
ret.push_back({var, builtin});
@ -112,28 +114,31 @@ Function::ReferencedBuiltinVariables() const {
return ret;
}
Function::VariableBindings Function::ReferencedSamplerVariables() const {
return ReferencedSamplerVariablesImpl(ast::SamplerKind::kSampler);
}
Function::VariableBindings Function::ReferencedComparisonSamplerVariables()
Function::VariableBindings Function::TransitivelyReferencedSamplerVariables()
const {
return ReferencedSamplerVariablesImpl(ast::SamplerKind::kComparisonSampler);
return TransitivelyReferencedSamplerVariablesImpl(ast::SamplerKind::kSampler);
}
Function::VariableBindings Function::ReferencedSampledTextureVariables() const {
return ReferencedSampledTextureVariablesImpl(false);
Function::VariableBindings
Function::TransitivelyReferencedComparisonSamplerVariables() const {
return TransitivelyReferencedSamplerVariablesImpl(
ast::SamplerKind::kComparisonSampler);
}
Function::VariableBindings Function::ReferencedMultisampledTextureVariables()
const {
return ReferencedSampledTextureVariablesImpl(true);
Function::VariableBindings
Function::TransitivelyReferencedSampledTextureVariables() const {
return TransitivelyReferencedSampledTextureVariablesImpl(false);
}
Function::VariableBindings Function::ReferencedVariablesOfType(
Function::VariableBindings
Function::TransitivelyReferencedMultisampledTextureVariables() const {
return TransitivelyReferencedSampledTextureVariablesImpl(true);
}
Function::VariableBindings Function::TransitivelyReferencedVariablesOfType(
const tint::TypeInfo& type_info) const {
VariableBindings ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
auto* unwrapped_type = var->Type()->UnwrapRef();
if (unwrapped_type->TypeInfo().Is(type_info)) {
if (auto binding_point = var->Declaration()->BindingPoint()) {
@ -153,11 +158,11 @@ bool Function::HasAncestorEntryPoint(Symbol symbol) const {
return false;
}
Function::VariableBindings Function::ReferencedSamplerVariablesImpl(
Function::VariableBindings Function::TransitivelyReferencedSamplerVariablesImpl(
ast::SamplerKind kind) const {
VariableBindings ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
auto* unwrapped_type = var->Type()->UnwrapRef();
auto* sampler = unwrapped_type->As<sem::Sampler>();
if (sampler == nullptr || sampler->kind() != kind) {
@ -171,11 +176,12 @@ Function::VariableBindings Function::ReferencedSamplerVariablesImpl(
return ret;
}
Function::VariableBindings Function::ReferencedSampledTextureVariablesImpl(
Function::VariableBindings
Function::TransitivelyReferencedSampledTextureVariablesImpl(
bool multisampled) const {
VariableBindings ret;
for (auto* var : ReferencedModuleVariables()) {
for (auto* var : TransitivelyReferencedGlobals()) {
auto* unwrapped_type = var->Type()->UnwrapRef();
auto* texture = unwrapped_type->As<sem::Texture>();
if (texture == nullptr) {

93
src/sem/function.h
View File

@ -20,23 +20,22 @@
#include <vector>
#include "src/ast/variable.h"
#include "src/sem/call_target.h"
#include "src/sem/call.h"
#include "src/utils/unique_vector.h"
namespace tint {
// Forward declarations
namespace ast {
class BindingDecoration;
class BuiltinDecoration;
class CallExpression;
class Function;
class GroupDecoration;
class LocationDecoration;
class ReturnStatement;
} // namespace ast
namespace sem {
class Intrinsic;
class Variable;
/// WorkgroupDimension describes the size of a single dimension of an entry
@ -49,6 +48,9 @@ struct WorkgroupDimension {
const ast::Variable* overridable_const = nullptr;
};
/// WorkgroupSize is a three-dimensional array of WorkgroupDimensions.
using WorkgroupSize = std::array<WorkgroupDimension, 3>;
/// Function holds the semantic information for function nodes.
class Function : public Castable<Function, CallTarget> {
public:
@ -60,21 +62,19 @@ class Function : public Castable<Function, CallTarget> {
/// @param declaration the ast::Function
/// @param return_type the return type of the function
/// @param parameters the parameters to the function
/// @param referenced_module_vars the referenced module variables
/// @param local_referenced_module_vars the locally referenced module
/// @param return_statements the function return statements
/// @param transitively_referenced_globals the referenced module variables
/// @param directly_referenced_globals the locally referenced module
/// @param callsites the callsites of the function
/// @param ancestor_entry_points the ancestor entry points
/// @param workgroup_size the workgroup size
Function(const ast::Function* declaration,
Type* return_type,
std::vector<Parameter*> parameters,
std::vector<const Variable*> referenced_module_vars,
std::vector<const Variable*> local_referenced_module_vars,
std::vector<const ast::ReturnStatement*> return_statements,
std::vector<const GlobalVariable*> transitively_referenced_globals,
std::vector<const GlobalVariable*> directly_referenced_globals,
std::vector<const ast::CallExpression*> callsites,
std::vector<Symbol> ancestor_entry_points,
std::array<WorkgroupDimension, 3> workgroup_size);
sem::WorkgroupSize workgroup_size);
/// Destructor
~Function() override;
@ -82,81 +82,78 @@ class Function : public Castable<Function, CallTarget> {
/// @returns the ast::Function declaration
const ast::Function* Declaration() const { return declaration_; }
/// 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<const Variable*>& ReferencedModuleVariables() const {
return referenced_module_vars_;
}
/// @returns the locally referenced module variables
const std::vector<const Variable*>& LocalReferencedModuleVariables() const {
return local_referenced_module_vars_;
}
/// @returns the return statements
const std::vector<const ast::ReturnStatement*> ReturnStatements() const {
return return_statements_;
/// @returns the workgroup size {x, y, z} for the function.
const sem::WorkgroupSize& WorkgroupSize() const { return workgroup_size_; }
/// @returns all transitively referenced global variables
const utils::UniqueVector<const GlobalVariable*>&
TransitivelyReferencedGlobals() const {
return transitively_referenced_globals_;
}
/// @returns the list of callsites of this function
std::vector<const ast::CallExpression*> CallSites() const {
return callsites_;
}
/// @returns the ancestor entry points
/// @returns the names of the ancestor entry points
const std::vector<Symbol>& AncestorEntryPoints() const {
return ancestor_entry_points_;
}
/// Retrieves any referenced location variables
/// @returns the <variable, decoration> pair.
std::vector<std::pair<const Variable*, const ast::LocationDecoration*>>
ReferencedLocationVariables() const;
TransitivelyReferencedLocationVariables() const;
/// Retrieves any referenced builtin variables
/// @returns the <variable, decoration> pair.
std::vector<std::pair<const Variable*, const ast::BuiltinDecoration*>>
ReferencedBuiltinVariables() const;
TransitivelyReferencedBuiltinVariables() const;
/// Retrieves any referenced uniform variables. Note, the variables must be
/// decorated with both binding and group decorations.
/// @returns the referenced uniforms
VariableBindings ReferencedUniformVariables() const;
VariableBindings TransitivelyReferencedUniformVariables() const;
/// Retrieves any referenced storagebuffer variables. Note, the variables
/// must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
VariableBindings ReferencedStorageBufferVariables() const;
VariableBindings TransitivelyReferencedStorageBufferVariables() const;
/// Retrieves any referenced regular Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
VariableBindings ReferencedSamplerVariables() const;
VariableBindings TransitivelyReferencedSamplerVariables() const;
/// Retrieves any referenced comparison Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
VariableBindings ReferencedComparisonSamplerVariables() const;
VariableBindings TransitivelyReferencedComparisonSamplerVariables() const;
/// Retrieves any referenced sampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
VariableBindings ReferencedSampledTextureVariables() const;
VariableBindings TransitivelyReferencedSampledTextureVariables() const;
/// Retrieves any referenced multisampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
VariableBindings ReferencedMultisampledTextureVariables() const;
VariableBindings TransitivelyReferencedMultisampledTextureVariables() const;
/// Retrieves any referenced variables of the given type. Note, the variables
/// must be decorated with both binding and group decorations.
/// @param type_info the type of the variables to find
/// @returns the referenced variables
VariableBindings ReferencedVariablesOfType(
VariableBindings TransitivelyReferencedVariablesOfType(
const tint::TypeInfo& type_info) const;
/// Retrieves any referenced variables of the given type. Note, the variables
/// must be decorated with both binding and group decorations.
/// @returns the referenced variables
template <typename T>
VariableBindings ReferencedVariablesOfType() const {
return ReferencedVariablesOfType(TypeInfo::Of<T>());
VariableBindings TransitivelyReferencedVariablesOfType() const {
return TransitivelyReferencedVariablesOfType(TypeInfo::Of<T>());
}
/// Checks if the given entry point is an ancestor
@ -164,23 +161,21 @@ class Function : public Castable<Function, CallTarget> {
/// @returns true if `sym` is an ancestor entry point of this function
bool HasAncestorEntryPoint(Symbol sym) const;
/// @returns the workgroup size {x, y, z} for the function.
const std::array<WorkgroupDimension, 3>& workgroup_size() const {
return workgroup_size_;
}
private:
VariableBindings ReferencedSamplerVariablesImpl(ast::SamplerKind kind) const;
VariableBindings ReferencedSampledTextureVariablesImpl(
VariableBindings TransitivelyReferencedSamplerVariablesImpl(
ast::SamplerKind kind) const;
VariableBindings TransitivelyReferencedSampledTextureVariablesImpl(
bool multisampled) const;
const ast::Function* const declaration_;
std::vector<const Variable*> const referenced_module_vars_;
std::vector<const Variable*> const local_referenced_module_vars_;
std::vector<const ast::ReturnStatement*> const return_statements_;
std::vector<const ast::CallExpression*> const callsites_;
std::vector<Symbol> const ancestor_entry_points_;
std::array<WorkgroupDimension, 3> workgroup_size_;
const sem::WorkgroupSize workgroup_size_;
utils::UniqueVector<const GlobalVariable*> directly_referenced_globals_;
utils::UniqueVector<const GlobalVariable*> transitively_referenced_globals_;
utils::UniqueVector<const Function*> transitively_called_functions_;
utils::UniqueVector<const Intrinsic*> directly_called_intrinsics_;
std::vector<const ast::CallExpression*> callsites_;
std::vector<Symbol> ancestor_entry_points_;
};
} // namespace sem

2
src/transform/binding_remapper.cc
View File

@ -64,7 +64,7 @@ void BindingRemapper::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) {
}
auto* func = ctx.src->Sem().Get(func_ast);
std::unordered_map<sem::BindingPoint, int> binding_point_counts;
for (auto* var : func->ReferencedModuleVariables()) {
for (auto* var : func->TransitivelyReferencedGlobals()) {
if (auto binding_point = var->Declaration()->BindingPoint()) {
BindingPoint from{binding_point.group->value,
binding_point.binding->value};

6
src/transform/module_scope_var_to_entry_point_param.cc
View File

@ -98,7 +98,7 @@ struct ModuleScopeVarToEntryPointParam::State {
auto* func_sem = ctx.src->Sem().Get(func_ast);
bool needs_processing = false;
for (auto* var : func_sem->ReferencedModuleVariables()) {
for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
if (var->StorageClass() != ast::StorageClass::kNone) {
needs_processing = true;
break;
@ -155,7 +155,7 @@ struct ModuleScopeVarToEntryPointParam::State {
return workgroup_parameter_symbol;
};
for (auto* var : func_sem->ReferencedModuleVariables()) {
for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
auto sc = var->StorageClass();
if (sc == ast::StorageClass::kNone) {
continue;
@ -312,7 +312,7 @@ struct ModuleScopeVarToEntryPointParam::State {
// Add new arguments for any variables that are needed by the callee.
// For entry points, pass non-handle types as pointers.
for (auto* target_var : target_sem->ReferencedModuleVariables()) {
for (auto* target_var : target_sem->TransitivelyReferencedGlobals()) {
auto sc = target_var->StorageClass();
if (sc == ast::StorageClass::kNone) {
continue;

2
src/transform/single_entry_point.cc
View File

@ -63,7 +63,7 @@ void SingleEntryPoint::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) {
// Build set of referenced module-scope variables for faster lookups later.
std::unordered_set<const ast::Variable*> referenced_vars;
for (auto* var : sem.Get(entry_point)->ReferencedModuleVariables()) {
for (auto* var : sem.Get(entry_point)->TransitivelyReferencedGlobals()) {
referenced_vars.emplace(var->Declaration());
}

2
src/transform/zero_init_workgroup_memory.cc
View File

@ -121,7 +121,7 @@ struct ZeroInitWorkgroupMemory::State {
// Generate a list of statements to zero initialize each of the
// workgroup storage variables used by `fn`. This will populate #statements.
auto* func = sem.Get(fn);
for (auto* var : func->ReferencedModuleVariables()) {
for (auto* var : func->TransitivelyReferencedGlobals()) {
if (var->StorageClass() == ast::StorageClass::kWorkgroup) {
BuildZeroingStatements(
var->Type()->UnwrapRef(), [&](uint32_t num_values) {

2
src/writer/glsl/generator_impl.cc
View File

@ -1756,7 +1756,7 @@ bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) {
auto out = line();
if (func->PipelineStage() == ast::PipelineStage::kCompute) {
// Emit the layout(local_size) attributes.
auto wgsize = func_sem->workgroup_size();
auto wgsize = func_sem->WorkgroupSize();
out << "layout(";
for (int i = 0; i < 3; i++) {
if (i > 0) {

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

@ -2580,7 +2580,7 @@ bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) {
auto out = line();
if (func->PipelineStage() == ast::PipelineStage::kCompute) {
// Emit the workgroup_size attribute.
auto wgsize = func_sem->workgroup_size();
auto wgsize = func_sem->WorkgroupSize();
out << "[numthreads(";
for (int i = 0; i < 3; i++) {
if (i > 0) {

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

@ -461,7 +461,7 @@ bool Builder::GenerateEntryPoint(const ast::Function* func, uint32_t id) {
Operand::String(builder_.Symbols().NameFor(func->symbol))};
auto* func_sem = builder_.Sem().Get(func);
for (const auto* var : func_sem->ReferencedModuleVariables()) {
for (const auto* var : func_sem->TransitivelyReferencedGlobals()) {
// 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.
if (var->StorageClass() != ast::StorageClass::kInput &&
@ -492,7 +492,7 @@ bool Builder::GenerateExecutionModes(const ast::Function* func, uint32_t id) {
spv::Op::OpExecutionMode,
{Operand::Int(id), Operand::Int(SpvExecutionModeOriginUpperLeft)});
} else if (func->PipelineStage() == ast::PipelineStage::kCompute) {
auto& wgsize = func_sem->workgroup_size();
auto& wgsize = func_sem->WorkgroupSize();
// Check if the workgroup_size uses pipeline-overridable constants.
if (wgsize[0].overridable_const || wgsize[1].overridable_const ||
@ -553,7 +553,7 @@ bool Builder::GenerateExecutionModes(const ast::Function* func, uint32_t id) {
}
}
for (auto builtin : func_sem->ReferencedBuiltinVariables()) {
for (auto builtin : func_sem->TransitivelyReferencedBuiltinVariables()) {
if (builtin.second->builtin == ast::Builtin::kFragDepth) {
push_execution_mode(
spv::Op::OpExecutionMode,