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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

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. // Build set of referenced module-scope variables for faster lookups later.
std::unordered_set<const ast::Variable*> referenced_vars; 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()); referenced_vars.emplace(var->Declaration());
} }

View File

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

View File

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

View File

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

View File

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