From b17aea159c328764238f0550d327bb531cf9bcf8 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 3 Feb 2021 17:51:09 +0000 Subject: [PATCH] Add semantic::Variable, use it. Pull the mutable semantic field from ast::Variable and into a new semantic::Variable node. Have the TypeDeterminer create these semantic::Variable nodes. Bug: tint:390 Change-Id: Ia13f5e7b065941ed66ea5a86c6ccb288071feff3 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/40063 Commit-Queue: Ben Clayton Reviewed-by: David Neto --- BUILD.gn | 1 + src/CMakeLists.txt | 1 + src/ast/variable.cc | 17 +-- src/ast/variable.h | 17 ++- src/ast/variable_test.cc | 6 +- src/inspector/inspector.cc | 64 +++++----- .../parser_impl_global_variable_decl_test.cc | 8 +- src/semantic/function.h | 37 +++--- src/semantic/sem_function.cc | 63 +++++----- src/semantic/sem_variable.cc | 28 +++++ src/semantic/type_mappings.h | 3 + src/semantic/variable.h | 64 ++++++++++ src/transform/emit_vertex_point_size.cc | 1 - src/transform/first_index_offset.cc | 4 +- src/transform/vertex_pulling.cc | 10 +- src/type_determiner.cc | 108 ++++++++++------ src/type_determiner.h | 62 ++++++---- src/type_determiner_test.cc | 25 ++-- src/validator/validator_impl.cc | 14 ++- src/writer/hlsl/generator_impl.cc | 79 ++++++------ src/writer/hlsl/generator_impl.h | 4 +- src/writer/msl/generator_impl.cc | 116 +++++++++--------- src/writer/msl/generator_impl.h | 6 +- src/writer/spirv/builder.cc | 21 ++-- src/writer/wgsl/generator_impl.cc | 19 +-- src/writer/wgsl/generator_impl.h | 2 +- 26 files changed, 485 insertions(+), 295 deletions(-) create mode 100644 src/semantic/sem_variable.cc create mode 100644 src/semantic/variable.h diff --git a/BUILD.gn b/BUILD.gn index 8e15afc8b5..12578bb58e 100644 --- a/BUILD.gn +++ b/BUILD.gn @@ -385,6 +385,7 @@ source_set("libtint_core_src") { "src/semantic/sem_function.cc", "src/semantic/sem_info.cc", "src/semantic/sem_node.cc", + "src/semantic/sem_variable.cc", "src/semantic/type_mappings.h", "src/source.cc", "src/source.h", diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 56a1efb9da..f79ae1a891 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -199,6 +199,7 @@ set(TINT_LIB_SRCS semantic/sem_function.cc semantic/sem_info.cc semantic/sem_node.cc + semantic/sem_variable.cc semantic/type_mappings.h source.cc source.h diff --git a/src/ast/variable.cc b/src/ast/variable.cc index cc0e676491..40056a4fb6 100644 --- a/src/ast/variable.cc +++ b/src/ast/variable.cc @@ -19,6 +19,7 @@ #include "src/ast/constant_id_decoration.h" #include "src/clone_context.h" #include "src/program_builder.h" +#include "src/semantic/variable.h" TINT_INSTANTIATE_CLASS_ID(tint::ast::Variable); @@ -38,7 +39,7 @@ Variable::Variable(const Source& source, is_const_(is_const), constructor_(constructor), decorations_(std::move(decorations)), - storage_class_(sc) {} + declared_storage_class_(sc) {} Variable::Variable(Variable&&) = default; @@ -91,10 +92,10 @@ uint32_t Variable::constant_id() const { } Variable* Variable::Clone(CloneContext* ctx) const { - return ctx->dst->create(ctx->Clone(source()), ctx->Clone(symbol_), - storage_class(), ctx->Clone(type()), - is_const_, ctx->Clone(constructor()), - ctx->Clone(decorations_)); + return ctx->dst->create( + ctx->Clone(source()), ctx->Clone(symbol_), declared_storage_class(), + ctx->Clone(type()), is_const_, ctx->Clone(constructor()), + ctx->Clone(decorations_)); } bool Variable::IsValid() const { @@ -110,13 +111,15 @@ bool Variable::IsValid() const { return true; } -void Variable::info_to_str(const semantic::Info&, +void Variable::info_to_str(const semantic::Info& sem, std::ostream& out, size_t indent) const { + auto* var_sem = sem.Get(this); make_indent(out, indent); out << symbol_.to_str() << std::endl; make_indent(out, indent); - out << storage_class_ << std::endl; + out << (var_sem ? var_sem->StorageClass() : declared_storage_class()) + << std::endl; make_indent(out, indent); out << type_->type_name() << std::endl; } diff --git a/src/ast/variable.h b/src/ast/variable.h index fb337b1f56..2d18b14ebf 100644 --- a/src/ast/variable.h +++ b/src/ast/variable.h @@ -84,7 +84,7 @@ class Variable : public Castable { /// Create a variable /// @param source the variable source /// @param sym the variable symbol - /// @param sc the variable storage class + /// @param sc the declared storage class /// @param type the value type /// @param is_const true if the variable is const /// @param constructor the constructor expression @@ -107,12 +107,10 @@ class Variable : public Castable { /// @returns the variable's type. type::Type* type() const { return type_; } - /// Sets the storage class - /// @param sc the storage class - void set_storage_class(StorageClass sc) { storage_class_ = sc; } - /// @returns the storage class - StorageClass storage_class() const { return storage_class_; } - + /// @returns the declared storage class + StorageClass declared_storage_class() const { + return declared_storage_class_; + } /// @returns the constructor expression or nullptr if none set Expression* constructor() const { return constructor_; } /// @returns true if the variable has an constructor @@ -126,7 +124,7 @@ class Variable : public Castable { /// @returns true if the decorations include a LocationDecoration bool HasLocationDecoration() const; - /// @returns true if the deocrations include a BuiltinDecoration + /// @returns true if the decorations include a BuiltinDecoration bool HasBuiltinDecoration() const; /// @returns true if the decorations include a ConstantIdDecoration bool HasConstantIdDecoration() const; @@ -182,8 +180,7 @@ class Variable : public Castable { bool const is_const_; Expression* const constructor_; VariableDecorationList const decorations_; - - StorageClass storage_class_ = StorageClass::kNone; // Semantic info + StorageClass const declared_storage_class_; }; /// A list of variables diff --git a/src/ast/variable_test.cc b/src/ast/variable_test.cc index f700381d57..9a9f6be27b 100644 --- a/src/ast/variable_test.cc +++ b/src/ast/variable_test.cc @@ -30,7 +30,7 @@ TEST_F(VariableTest, Creation) { auto* v = Var("my_var", StorageClass::kFunction, ty.i32()); EXPECT_EQ(v->symbol(), Symbol(1)); - EXPECT_EQ(v->storage_class(), StorageClass::kFunction); + EXPECT_EQ(v->declared_storage_class(), StorageClass::kFunction); EXPECT_EQ(v->type(), ty.i32()); EXPECT_EQ(v->source().range.begin.line, 0u); EXPECT_EQ(v->source().range.begin.column, 0u); @@ -44,7 +44,7 @@ TEST_F(VariableTest, CreationWithSource) { "i", StorageClass::kPrivate, ty.f32(), nullptr, VariableDecorationList{}); EXPECT_EQ(v->symbol(), Symbol(1)); - EXPECT_EQ(v->storage_class(), StorageClass::kPrivate); + EXPECT_EQ(v->declared_storage_class(), StorageClass::kPrivate); EXPECT_EQ(v->type(), ty.f32()); EXPECT_EQ(v->source().range.begin.line, 27u); EXPECT_EQ(v->source().range.begin.column, 4u); @@ -59,7 +59,7 @@ TEST_F(VariableTest, CreationEmpty) { VariableDecorationList{}); EXPECT_EQ(v->symbol(), Symbol(1)); - EXPECT_EQ(v->storage_class(), StorageClass::kWorkgroup); + EXPECT_EQ(v->declared_storage_class(), StorageClass::kWorkgroup); EXPECT_EQ(v->type(), ty.i32()); EXPECT_EQ(v->source().range.begin.line, 27u); EXPECT_EQ(v->source().range.begin.column, 4u); diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc index ccfe471bc9..ca4bb9767f 100644 --- a/src/inspector/inspector.cc +++ b/src/inspector/inspector.cc @@ -30,6 +30,7 @@ #include "src/ast/variable.h" #include "src/program.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/access_control_type.h" #include "src/type/array_type.h" #include "src/type/f32_type.h" @@ -66,14 +67,16 @@ std::vector Inspector::GetEntryPoints() { entry_point.workgroup_size_z) = func->workgroup_size(); for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) { - auto name = program_->Symbols().NameFor(var->symbol()); - if (var->HasBuiltinDecoration()) { + auto* decl = var->Declaration(); + + auto name = program_->Symbols().NameFor(decl->symbol()); + if (decl->HasBuiltinDecoration()) { continue; } StageVariable stage_variable; stage_variable.name = name; - auto* location_decoration = var->GetLocationDecoration(); + auto* location_decoration = decl->GetLocationDecoration(); if (location_decoration) { stage_variable.has_location_decoration = true; stage_variable.location_decoration = location_decoration->value(); @@ -81,9 +84,9 @@ std::vector Inspector::GetEntryPoints() { stage_variable.has_location_decoration = false; } - if (var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kInput) { entry_point.input_variables.push_back(stage_variable); - } else if (var->storage_class() == ast::StorageClass::kOutput) { + } else if (var->StorageClass() == ast::StorageClass::kOutput) { entry_point.output_variables.push_back(stage_variable); } } @@ -188,14 +191,14 @@ std::vector Inspector::GetUniformBufferResourceBindings( auto* func_sem = program_->Sem().Get(func); for (auto& ruv : func_sem->ReferencedUniformVariables()) { - ResourceBinding entry; - ast::Variable* var = nullptr; - semantic::Function::BindingInfo binding_info; - std::tie(var, binding_info) = ruv; - if (!var->type()->Is()) { + auto* var = ruv.first; + auto* decl = var->Declaration(); + auto binding_info = ruv.second; + + if (!decl->type()->Is()) { continue; } - auto* unwrapped_type = var->type()->UnwrapIfNeeded(); + auto* unwrapped_type = decl->type()->UnwrapIfNeeded(); auto* str = unwrapped_type->As(); if (str == nullptr) { @@ -206,10 +209,11 @@ std::vector Inspector::GetUniformBufferResourceBindings( continue; } + ResourceBinding entry; entry.bind_group = binding_info.group->value(); entry.binding = binding_info.binding->value(); entry.min_buffer_binding_size = - var->type()->MinBufferBindingSize(type::MemoryLayout::kUniformBuffer); + decl->type()->MinBufferBindingSize(type::MemoryLayout::kUniformBuffer); result.push_back(entry); } @@ -239,11 +243,9 @@ std::vector Inspector::GetSamplerResourceBindings( auto* func_sem = program_->Sem().Get(func); for (auto& rs : func_sem->ReferencedSamplerVariables()) { - ResourceBinding entry; - ast::Variable* var = nullptr; - semantic::Function::BindingInfo binding_info; - std::tie(var, binding_info) = rs; + auto binding_info = rs.second; + ResourceBinding entry; entry.bind_group = binding_info.group->value(); entry.binding = binding_info.binding->value(); @@ -264,11 +266,9 @@ std::vector Inspector::GetComparisonSamplerResourceBindings( auto* func_sem = program_->Sem().Get(func); for (auto& rcs : func_sem->ReferencedComparisonSamplerVariables()) { - ResourceBinding entry; - ast::Variable* var = nullptr; - semantic::Function::BindingInfo binding_info; - std::tie(var, binding_info) = rcs; + auto binding_info = rcs.second; + ResourceBinding entry; entry.bind_group = binding_info.group->value(); entry.binding = binding_info.binding->value(); @@ -314,12 +314,11 @@ std::vector Inspector::GetStorageBufferResourceBindingsImpl( auto* func_sem = program_->Sem().Get(func); std::vector result; for (auto& rsv : func_sem->ReferencedStoragebufferVariables()) { - ResourceBinding entry; - ast::Variable* var = nullptr; - semantic::Function::BindingInfo binding_info; - std::tie(var, binding_info) = rsv; + auto* var = rsv.first; + auto* decl = var->Declaration(); + auto binding_info = rsv.second; - auto* ac_type = var->type()->As(); + auto* ac_type = decl->type()->As(); if (ac_type == nullptr) { continue; } @@ -328,14 +327,15 @@ std::vector Inspector::GetStorageBufferResourceBindingsImpl( continue; } - if (!var->type()->UnwrapIfNeeded()->Is()) { + if (!decl->type()->UnwrapIfNeeded()->Is()) { continue; } + ResourceBinding entry; entry.bind_group = binding_info.group->value(); entry.binding = binding_info.binding->value(); entry.min_buffer_binding_size = - var->type()->MinBufferBindingSize(type::MemoryLayout::kStorageBuffer); + decl->type()->MinBufferBindingSize(type::MemoryLayout::kStorageBuffer); result.push_back(entry); } @@ -357,15 +357,15 @@ std::vector Inspector::GetSampledTextureResourceBindingsImpl( multisampled_only ? func_sem->ReferencedMultisampledTextureVariables() : func_sem->ReferencedSampledTextureVariables(); for (auto& ref : referenced_variables) { - ResourceBinding entry; - ast::Variable* var = nullptr; - semantic::Function::BindingInfo binding_info; - std::tie(var, binding_info) = ref; + auto* var = ref.first; + auto* decl = var->Declaration(); + auto binding_info = ref.second; + ResourceBinding entry; entry.bind_group = binding_info.group->value(); entry.binding = binding_info.binding->value(); - auto* texture_type = var->type()->UnwrapIfNeeded()->As(); + auto* texture_type = decl->type()->UnwrapIfNeeded()->As(); switch (texture_type->dim()) { case type::TextureDimension::k1d: entry.dim = ResourceBinding::TextureDimension::k1d; diff --git a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc index 9c83e4221a..625abb59b4 100644 --- a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc +++ b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc @@ -38,7 +38,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithoutConstructor) { EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a")); EXPECT_TRUE(e->type()->Is()); - EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput); + EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput); EXPECT_EQ(e->source().range.begin.line, 1u); EXPECT_EQ(e->source().range.begin.column, 10u); @@ -61,7 +61,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithConstructor) { EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a")); EXPECT_TRUE(e->type()->Is()); - EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput); + EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput); EXPECT_EQ(e->source().range.begin.line, 1u); EXPECT_EQ(e->source().range.begin.column, 10u); @@ -87,7 +87,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration) { EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a")); ASSERT_NE(e->type(), nullptr); EXPECT_TRUE(e->type()->Is()); - EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput); + EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput); EXPECT_EQ(e->source().range.begin.line, 1u); EXPECT_EQ(e->source().range.begin.column, 35u); @@ -117,7 +117,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration_MulitpleGroups) { EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a")); ASSERT_NE(e->type(), nullptr); EXPECT_TRUE(e->type()->Is()); - EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput); + EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput); EXPECT_EQ(e->source().range.begin.line, 1u); EXPECT_EQ(e->source().range.begin.column, 38u); diff --git a/src/semantic/function.h b/src/semantic/function.h index 0c751138b9..70708d986c 100644 --- a/src/semantic/function.h +++ b/src/semantic/function.h @@ -28,7 +28,6 @@ namespace tint { namespace ast { class BindingDecoration; class GroupDecoration; -class Variable; class LocationDecoration; class BuiltinDecoration; } // namespace ast @@ -38,6 +37,8 @@ class Type; namespace semantic { +class Variable; + /// Function holds the semantic information for function nodes. class Function : public Castable { public: @@ -54,8 +55,8 @@ class Function : public Castable { /// @param local_referenced_module_vars the locally referenced module /// variables /// @param ancestor_entry_points the ancestor entry points - explicit Function(std::vector referenced_module_vars, - std::vector local_referenced_module_vars, + explicit Function(std::vector referenced_module_vars, + std::vector local_referenced_module_vars, std::vector ancestor_entry_points); /// Destructor @@ -64,11 +65,11 @@ class Function : public Castable { /// 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& ReferencedModuleVariables() const { + const std::vector& ReferencedModuleVariables() const { return referenced_module_vars_; } /// @returns the locally referenced module variables - const std::vector& LocalReferencedModuleVariables() const { + const std::vector& LocalReferencedModuleVariables() const { return local_referenced_module_vars_; } /// @returns the ancestor entry points @@ -77,53 +78,53 @@ class Function : public Castable { } /// Retrieves any referenced location variables /// @returns the pair. - const std::vector> + const std::vector> ReferencedLocationVariables() const; /// Retrieves any referenced builtin variables /// @returns the pair. - const std::vector> + const std::vector> 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> + const std::vector> 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> + const std::vector> 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> + const std::vector> 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> + const std::vector> 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> + const std::vector> 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> + const std::vector> ReferencedMultisampledTextureVariables() const; /// Retrieves any locally referenced builtin variables /// @returns the pairs. - const std::vector> + const std::vector> LocalReferencedBuiltinVariables() const; /// Checks if the given entry point is an ancestor @@ -132,13 +133,13 @@ class Function : public Castable { bool HasAncestorEntryPoint(Symbol sym) const; private: - const std::vector> + const std::vector> ReferencedSamplerVariablesImpl(type::SamplerKind kind) const; - const std::vector> + const std::vector> ReferencedSampledTextureVariablesImpl(bool multisampled) const; - std::vector const referenced_module_vars_; - std::vector const local_referenced_module_vars_; + std::vector const referenced_module_vars_; + std::vector const local_referenced_module_vars_; std::vector const ancestor_entry_points_; }; diff --git a/src/semantic/sem_function.cc b/src/semantic/sem_function.cc index 561369702f..26768ec9fc 100644 --- a/src/semantic/sem_function.cc +++ b/src/semantic/sem_function.cc @@ -20,6 +20,7 @@ #include "src/ast/location_decoration.h" #include "src/ast/variable.h" #include "src/ast/variable_decoration.h" +#include "src/semantic/variable.h" #include "src/type/multisampled_texture_type.h" #include "src/type/sampled_texture_type.h" #include "src/type/texture_type.h" @@ -29,8 +30,8 @@ TINT_INSTANTIATE_CLASS_ID(tint::semantic::Function); namespace tint { namespace semantic { -Function::Function(std::vector referenced_module_vars, - std::vector local_referenced_module_vars, +Function::Function(std::vector referenced_module_vars, + std::vector local_referenced_module_vars, std::vector ancestor_entry_points) : referenced_module_vars_(std::move(referenced_module_vars)), local_referenced_module_vars_(std::move(local_referenced_module_vars)), @@ -38,12 +39,12 @@ Function::Function(std::vector referenced_module_vars, Function::~Function() = default; -const std::vector> +const std::vector> Function::ReferencedLocationVariables() const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* location = deco->As()) { ret.push_back({var, location}); break; @@ -53,18 +54,18 @@ Function::ReferencedLocationVariables() const { return ret; } -const std::vector> +const std::vector> Function::ReferencedUniformVariables() const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - if (var->storage_class() != ast::StorageClass::kUniform) { + if (var->StorageClass() != ast::StorageClass::kUniform) { continue; } ast::BindingDecoration* binding = nullptr; ast::GroupDecoration* group = nullptr; - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* b = deco->As()) { binding = b; } else if (auto* g = deco->As()) { @@ -80,18 +81,18 @@ Function::ReferencedUniformVariables() const { return ret; } -const std::vector> +const std::vector> Function::ReferencedStoragebufferVariables() const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - if (var->storage_class() != ast::StorageClass::kStorage) { + if (var->StorageClass() != ast::StorageClass::kStorage) { continue; } ast::BindingDecoration* binding = nullptr; ast::GroupDecoration* group = nullptr; - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* b = deco->As()) { binding = b; } else if (auto* s = deco->As()) { @@ -107,12 +108,12 @@ Function::ReferencedStoragebufferVariables() const { return ret; } -const std::vector> +const std::vector> Function::ReferencedBuiltinVariables() const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* builtin = deco->As()) { ret.push_back({var, builtin}); break; @@ -122,32 +123,32 @@ Function::ReferencedBuiltinVariables() const { return ret; } -const std::vector> +const std::vector> Function::ReferencedSamplerVariables() const { return ReferencedSamplerVariablesImpl(type::SamplerKind::kSampler); } -const std::vector> +const std::vector> Function::ReferencedComparisonSamplerVariables() const { return ReferencedSamplerVariablesImpl(type::SamplerKind::kComparisonSampler); } -const std::vector> +const std::vector> Function::ReferencedSampledTextureVariables() const { return ReferencedSampledTextureVariablesImpl(false); } -const std::vector> +const std::vector> Function::ReferencedMultisampledTextureVariables() const { return ReferencedSampledTextureVariablesImpl(true); } -const std::vector> +const std::vector> Function::LocalReferencedBuiltinVariables() const { - std::vector> ret; + std::vector> ret; for (auto* var : LocalReferencedModuleVariables()) { - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* builtin = deco->As()) { ret.push_back({var, builtin}); break; @@ -166,12 +167,12 @@ bool Function::HasAncestorEntryPoint(Symbol symbol) const { return false; } -const std::vector> +const std::vector> Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - auto* unwrapped_type = var->type()->UnwrapIfNeeded(); + auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded(); auto* sampler = unwrapped_type->As(); if (sampler == nullptr || sampler->kind() != kind) { continue; @@ -179,7 +180,7 @@ Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const { ast::BindingDecoration* binding = nullptr; ast::GroupDecoration* group = nullptr; - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* b = deco->As()) { binding = b; } @@ -196,12 +197,12 @@ Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const { return ret; } -const std::vector> +const std::vector> Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const { - std::vector> ret; + std::vector> ret; for (auto* var : ReferencedModuleVariables()) { - auto* unwrapped_type = var->type()->UnwrapIfNeeded(); + auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded(); auto* texture = unwrapped_type->As(); if (texture == nullptr) { continue; @@ -216,7 +217,7 @@ Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const { ast::BindingDecoration* binding = nullptr; ast::GroupDecoration* group = nullptr; - for (auto* deco : var->decorations()) { + for (auto* deco : var->Declaration()->decorations()) { if (auto* b = deco->As()) { binding = b; } else if (auto* s = deco->As()) { diff --git a/src/semantic/sem_variable.cc b/src/semantic/sem_variable.cc new file mode 100644 index 0000000000..33e84a1d82 --- /dev/null +++ b/src/semantic/sem_variable.cc @@ -0,0 +1,28 @@ +// 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/variable.h" + +TINT_INSTANTIATE_CLASS_ID(tint::semantic::Variable); + +namespace tint { +namespace semantic { + +Variable::Variable(ast::Variable* declaration, ast::StorageClass storage_class) + : declaration_(declaration), storage_class_(storage_class) {} + +Variable::~Variable() = default; + +} // namespace semantic +} // namespace tint diff --git a/src/semantic/type_mappings.h b/src/semantic/type_mappings.h index ac099a8c59..58b91c5428 100644 --- a/src/semantic/type_mappings.h +++ b/src/semantic/type_mappings.h @@ -24,6 +24,7 @@ namespace ast { class Expression; class Function; +class Variable; } // namespace ast @@ -31,6 +32,7 @@ namespace semantic { class Expression; class Function; +class Variable; /// TypeMappings is a struct that holds dummy `operator()` methods that's used /// by SemanticNodeTypeFor to map AST node types to their corresponding semantic @@ -41,6 +43,7 @@ struct TypeMappings { //! @cond Doxygen_Suppress semantic::Expression* operator()(ast::Expression*); semantic::Function* operator()(ast::Function*); + semantic::Variable* operator()(ast::Variable*); //! @endcond }; diff --git a/src/semantic/variable.h b/src/semantic/variable.h new file mode 100644 index 0000000000..6f15dc7506 --- /dev/null +++ b/src/semantic/variable.h @@ -0,0 +1,64 @@ +// 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_VARIABLE_H_ +#define SRC_SEMANTIC_VARIABLE_H_ + +#include +#include + +#include "src/ast/storage_class.h" +#include "src/semantic/node.h" +#include "src/type/sampler_type.h" + +namespace tint { + +// Forward declarations +namespace ast { +class Variable; +} // namespace ast +namespace type { +class Type; +} // namespace type + +namespace semantic { + +/// Variable holds the semantic information for variables. +class Variable : public Castable { + public: + /// Constructor + /// @param declaration the AST declaration node + /// @param storage_class the variable storage class + explicit Variable(ast::Variable* declaration, + ast::StorageClass storage_class); + + /// Destructor + ~Variable() override; + + /// @returns the AST declaration node + ast::Variable* Declaration() const { return declaration_; } + + /// @returns the storage class for the variable + ast::StorageClass StorageClass() const { return storage_class_; } + + private: + ast::Variable* const declaration_; + ast::StorageClass const storage_class_; +}; + +} // namespace semantic +} // namespace tint + +#endif // SRC_SEMANTIC_VARIABLE_H_ diff --git a/src/transform/emit_vertex_point_size.cc b/src/transform/emit_vertex_point_size.cc index 319d14434a..748021cd49 100644 --- a/src/transform/emit_vertex_point_size.cc +++ b/src/transform/emit_vertex_point_size.cc @@ -58,7 +58,6 @@ Transform::Output EmitVertexPointSize::Run(const Program* in) { false, // is_const nullptr, // constructor ast::VariableDecorationList{ - // decorations out.create(Source{}, ast::Builtin::kPointSize), }); diff --git a/src/transform/first_index_offset.cc b/src/transform/first_index_offset.cc index fec1068960..4ed52ab437 100644 --- a/src/transform/first_index_offset.cc +++ b/src/transform/first_index_offset.cc @@ -70,7 +70,7 @@ ast::Variable* clone_variable_with_new_name(CloneContext* ctx, return ctx->dst->create( ctx->Clone(in->source()), // source ctx->dst->Symbols().Register(new_name), // symbol - in->storage_class(), // storage_class + in->declared_storage_class(), // declared_storage_class ctx->Clone(in->type()), // type in->is_const(), // is_const ctx->Clone(in->constructor()), // constructor @@ -226,7 +226,7 @@ ast::Variable* FirstIndexOffset::AddUniformBuffer(ProgramBuilder* dst) { ast::VariableDecorationList{ dst->create(Source{}, binding_), dst->create(Source{}, group_), - }); // decorations + }); dst->AST().AddGlobalVariable(idx_var); diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc index 39585bea7f..f052bd44aa 100644 --- a/src/transform/vertex_pulling.cc +++ b/src/transform/vertex_pulling.cc @@ -35,6 +35,7 @@ #include "src/clone_context.h" #include "src/program.h" #include "src/program_builder.h" +#include "src/semantic/variable.h" #include "src/type/array_type.h" #include "src/type/f32_type.h" #include "src/type/i32_type.h" @@ -150,7 +151,8 @@ void VertexPulling::State::FindOrInsertVertexIndexIfUsed() { // Look for an existing vertex index builtin for (auto* v : ctx.src->AST().GlobalVariables()) { - if (v->storage_class() != ast::StorageClass::kInput) { + auto* sem = ctx.src->Sem().Get(v); + if (sem->StorageClass() != ast::StorageClass::kInput) { continue; } @@ -196,7 +198,8 @@ void VertexPulling::State::FindOrInsertInstanceIndexIfUsed() { // Look for an existing instance index builtin for (auto* v : ctx.src->AST().GlobalVariables()) { - if (v->storage_class() != ast::StorageClass::kInput) { + auto* sem = ctx.src->Sem().Get(v); + if (sem->StorageClass() != ast::StorageClass::kInput) { continue; } @@ -229,7 +232,8 @@ void VertexPulling::State::FindOrInsertInstanceIndexIfUsed() { void VertexPulling::State::ConvertVertexInputVariablesToPrivate() { for (auto* v : ctx.src->AST().GlobalVariables()) { - if (v->storage_class() != ast::StorageClass::kInput) { + auto* sem = ctx.src->Sem().Get(v); + if (sem->StorageClass() != ast::StorageClass::kInput) { continue; } diff --git a/src/type_determiner.cc b/src/type_determiner.cc index 61723843ab..e75bdc8f49 100644 --- a/src/type_determiner.cc +++ b/src/type_determiner.cc @@ -45,6 +45,7 @@ #include "src/program_builder.h" #include "src/semantic/expression.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/array_type.h" #include "src/type/bool_type.h" #include "src/type/depth_texture_type.h" @@ -88,13 +89,13 @@ void TypeDeterminer::set_error(const Source& src, const std::string& msg) { error_ += msg; } -void TypeDeterminer::set_referenced_from_function_if_needed(ast::Variable* var, +void TypeDeterminer::set_referenced_from_function_if_needed(VariableInfo* var, bool local) { if (current_function_ == nullptr) { return; } - if (var->storage_class() == ast::StorageClass::kNone || - var->storage_class() == ast::StorageClass::kFunction) { + if (var->storage_class == ast::StorageClass::kNone || + var->storage_class == ast::StorageClass::kFunction) { return; } @@ -105,7 +106,18 @@ void TypeDeterminer::set_referenced_from_function_if_needed(ast::Variable* var, } bool TypeDeterminer::Determine() { + bool result = DetermineInternal(); + + // Even if resolving failed, create all the semantic nodes for information we + // did generate. + CreateSemanticNodes(); + + return result; +} + +bool TypeDeterminer::DetermineInternal() { std::vector storage_textures; + for (auto& it : builder_->Types().types()) { if (auto* storage = it.second->UnwrapIfNeeded()->As()) { @@ -122,7 +134,7 @@ bool TypeDeterminer::Determine() { } for (auto* var : builder_->AST().GlobalVariables()) { - variable_stack_.set_global(var->symbol(), var); + variable_stack_.set_global(var->symbol(), CreateVariableInfo(var)); if (var->has_constructor()) { if (!DetermineResultType(var->constructor())) { @@ -135,8 +147,8 @@ bool TypeDeterminer::Determine() { return false; } - // Walk over the caller to callee information and update functions with which - // entry points call those functions. + // Walk over the caller to callee information and update functions with + // which entry points call those functions. for (auto* func : builder_->AST().Functions()) { if (!func->IsEntryPoint()) { continue; @@ -146,8 +158,6 @@ bool TypeDeterminer::Determine() { } } - CreateSemanticFunctions(); - return true; } @@ -170,15 +180,13 @@ bool TypeDeterminer::DetermineFunctions(const ast::FunctionList& funcs) { } bool TypeDeterminer::DetermineFunction(ast::Function* func) { - auto* info = function_infos_.Create(func); - symbol_to_function_[func->symbol()] = info; - function_to_info_.emplace(func, info); - - current_function_ = info; + current_function_ = function_infos_.Create(func); + symbol_to_function_[func->symbol()] = current_function_; + function_to_info_.emplace(func, current_function_); variable_stack_.push_scope(); for (auto* param : func->params()) { - variable_stack_.set(param->symbol(), param); + variable_stack_.set(param->symbol(), CreateVariableInfo(param)); } if (!DetermineStatements(func->body())) { @@ -211,22 +219,26 @@ bool TypeDeterminer::DetermineVariableStorageClass(ast::Statement* stmt) { } auto* var = var_decl->variable(); + + auto* info = CreateVariableInfo(var); + variable_to_info_.emplace(var, info); + // Nothing to do for const if (var->is_const()) { return true; } - if (var->storage_class() == ast::StorageClass::kFunction) { + if (info->storage_class == ast::StorageClass::kFunction) { return true; } - if (var->storage_class() != ast::StorageClass::kNone) { + if (info->storage_class != ast::StorageClass::kNone) { set_error(stmt->source(), "function variable has a non-function storage class"); return false; } - var->set_storage_class(ast::StorageClass::kFunction); + info->storage_class = ast::StorageClass::kFunction; return true; } @@ -291,7 +303,8 @@ bool TypeDeterminer::DetermineResultType(ast::Statement* stmt) { return true; } if (auto* v = stmt->As()) { - variable_stack_.set(v->variable()->symbol(), v->variable()); + variable_stack_.set(v->variable()->symbol(), + variable_to_info_.at(v->variable())); return DetermineResultType(v->variable()->constructor()); } @@ -816,17 +829,17 @@ bool TypeDeterminer::DetermineConstructor(ast::ConstructorExpression* expr) { bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) { auto symbol = expr->symbol(); - ast::Variable* var; + VariableInfo* var; if (variable_stack_.get(symbol, &var)) { // A constant is the type, but a variable is always a pointer so synthesize // the pointer around the variable type. - if (var->is_const()) { - SetType(expr, var->type()); - } else if (var->type()->Is()) { - SetType(expr, var->type()); + if (var->declaration->is_const()) { + SetType(expr, var->declaration->type()); + } else if (var->declaration->type()->Is()) { + SetType(expr, var->declaration->type()); } else { - SetType(expr, builder_->create(var->type(), - var->storage_class())); + SetType(expr, builder_->create(var->declaration->type(), + var->storage_class)); } set_referenced_from_function_if_needed(var, true); @@ -1146,6 +1159,13 @@ bool TypeDeterminer::DetermineUnaryOp(ast::UnaryOpExpression* expr) { return true; } +TypeDeterminer::VariableInfo* TypeDeterminer::CreateVariableInfo( + ast::Variable* var) { + auto* info = variable_infos_.Create(var); + variable_to_info_.emplace(var, info); + return info; +} + bool TypeDeterminer::DetermineStorageTextureSubtype(type::StorageTexture* tex) { if (tex->type() != nullptr) { return true; @@ -1211,22 +1231,40 @@ void TypeDeterminer::SetType(ast::Expression* expr, type::Type* type) const { builder_->create(type)); } -void TypeDeterminer::CreateSemanticFunctions() const { +void TypeDeterminer::CreateSemanticNodes() const { + auto& sem = builder_->Sem(); + + for (auto it : variable_to_info_) { + auto* var = it.first; + auto* info = it.second; + sem.Add(var, + builder_->create(var, info->storage_class)); + } + + auto remap_vars = [&sem](const std::vector& in) { + std::vector out; + out.reserve(in.size()); + for (auto* info : in) { + out.emplace_back(sem.Get(info->declaration)); + } + return out; + }; + 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( - info->referenced_module_vars, - info->local_referenced_module_vars, - info->ancestor_entry_points)); + sem.Add(func, builder_->create( + remap_vars(info->referenced_module_vars), + remap_vars(info->local_referenced_module_vars), + info->ancestor_entry_points)); } } +TypeDeterminer::VariableInfo::VariableInfo(ast::Variable* decl) + : declaration(decl), storage_class(decl->declared_storage_class()) {} + +TypeDeterminer::VariableInfo::~VariableInfo() = default; + TypeDeterminer::FunctionInfo::FunctionInfo(ast::Function* decl) : declaration(decl) {} diff --git a/src/type_determiner.h b/src/type_determiner.h index c943a536b1..1402dda36a 100644 --- a/src/type_determiner.h +++ b/src/type_determiner.h @@ -66,23 +66,6 @@ class TypeDeterminer { /// @returns true if the type determiner was successful bool Determine(); - /// Creates the semantic::Function nodes and adds them to the semantic::Info - /// of the ProgramBuilder. - void CreateSemanticFunctions() const; - - /// Retrieves information for the requested import. - /// @param src the source of the import - /// @param path the import path - /// @param name the method name to get information on - /// @param params the parameters to the method call - /// @param id out parameter for the external call ID. Must not be a nullptr. - /// @returns the return type of `name` in `path` or nullptr on error. - type::Type* GetImportData(const Source& src, - const std::string& path, - const std::string& name, - const ast::ExpressionList& params, - uint32_t* id); - /// Sets the intrinsic data information for the identifier if needed /// @param ident the identifier expression /// @returns true if an intrinsic was set @@ -99,6 +82,7 @@ class TypeDeterminer { set.emplace(val); } } + size_t size() const { return vector.size(); } ConstIterator begin() const { return vector.begin(); } ConstIterator end() const { return vector.end(); } operator const std::vector &() const { return vector; } @@ -108,6 +92,16 @@ class TypeDeterminer { std::unordered_set set; }; + /// Structure holding semantic information about a function. + /// Used to build the semantic::Function nodes at the end of resolving. + struct VariableInfo { + explicit VariableInfo(ast::Variable* decl); + ~VariableInfo(); + + ast::Variable* const declaration; + ast::StorageClass storage_class; + }; + /// Structure holding semantic information about a function. /// Used to build the semantic::Function nodes at the end of resolving. struct FunctionInfo { @@ -115,11 +109,16 @@ class TypeDeterminer { ~FunctionInfo(); ast::Function* const declaration; - UniqueVector referenced_module_vars; - UniqueVector local_referenced_module_vars; + UniqueVector referenced_module_vars; + UniqueVector local_referenced_module_vars; UniqueVector ancestor_entry_points; }; + /// Determines type information for the program, without creating final the + /// semantic nodes. + /// @returns true if the determination was successful + bool DetermineInternal(); + /// Determines type information for functions /// @param funcs the functions to check /// @returns true if the determination was successful @@ -154,8 +153,25 @@ class TypeDeterminer { /// @returns false on error bool DetermineStorageTextureSubtype(type::StorageTexture* tex); + /// Creates the nodes and adds them to the semantic::Info mappings of the + /// ProgramBuilder. + void CreateSemanticNodes() const; + + /// Retrieves information for the requested import. + /// @param src the source of the import + /// @param path the import path + /// @param name the method name to get information on + /// @param params the parameters to the method call + /// @param id out parameter for the external call ID. Must not be a nullptr. + /// @returns the return type of `name` in `path` or nullptr on error. + type::Type* GetImportData(const Source& src, + const std::string& path, + const std::string& name, + const ast::ExpressionList& params, + uint32_t* id); + 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(VariableInfo* var, bool local); void set_entry_points(const Symbol& fn_sym, Symbol ep_sym); bool DetermineArrayAccessor(ast::ArrayAccessorExpression* expr); @@ -169,6 +185,8 @@ class TypeDeterminer { bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr); bool DetermineUnaryOp(ast::UnaryOpExpression* expr); + VariableInfo* CreateVariableInfo(ast::Variable*); + /// @returns the resolved type of the ast::Expression `expr` /// @param expr the expression type::Type* TypeOf(ast::Expression* expr) const { @@ -183,10 +201,12 @@ class TypeDeterminer { ProgramBuilder* builder_; std::string error_; - ScopeStack variable_stack_; + ScopeStack variable_stack_; std::unordered_map symbol_to_function_; std::unordered_map function_to_info_; + std::unordered_map variable_to_info_; FunctionInfo* current_function_ = nullptr; + BlockAllocator variable_infos_; BlockAllocator function_infos_; // Map from caller functions to callee functions. diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc index b97f009ba4..96def378b2 100644 --- a/src/type_determiner_test.cc +++ b/src/type_determiner_test.cc @@ -53,6 +53,7 @@ #include "src/program_builder.h" #include "src/semantic/expression.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" #include "src/type/bool_type.h" @@ -689,11 +690,11 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables) { const auto& vars = func_sem->ReferencedModuleVariables(); ASSERT_EQ(vars.size(), 5u); - EXPECT_EQ(vars[0], out_var); - EXPECT_EQ(vars[1], in_var); - EXPECT_EQ(vars[2], wg_var); - EXPECT_EQ(vars[3], sb_var); - EXPECT_EQ(vars[4], priv_var); + EXPECT_EQ(vars[0]->Declaration(), out_var); + EXPECT_EQ(vars[1]->Declaration(), in_var); + EXPECT_EQ(vars[2]->Declaration(), wg_var); + EXPECT_EQ(vars[3]->Declaration(), sb_var); + EXPECT_EQ(vars[4]->Declaration(), priv_var); } TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) { @@ -726,11 +727,11 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) { const auto& vars = func2_sem->ReferencedModuleVariables(); ASSERT_EQ(vars.size(), 5u); - EXPECT_EQ(vars[0], out_var); - EXPECT_EQ(vars[1], in_var); - EXPECT_EQ(vars[2], wg_var); - EXPECT_EQ(vars[3], sb_var); - EXPECT_EQ(vars[4], priv_var); + EXPECT_EQ(vars[0]->Declaration(), out_var); + EXPECT_EQ(vars[1]->Declaration(), in_var); + EXPECT_EQ(vars[2]->Declaration(), wg_var); + EXPECT_EQ(vars[3]->Declaration(), sb_var); + EXPECT_EQ(vars[4]->Declaration(), priv_var); } TEST_F(TypeDeterminerTest, Function_NotRegisterFunctionVariable) { @@ -1490,7 +1491,7 @@ TEST_F(TypeDeterminerTest, StorageClass_SetsIfMissing) { EXPECT_TRUE(td()->Determine()) << td()->error(); - EXPECT_EQ(var->storage_class(), ast::StorageClass::kFunction); + EXPECT_EQ(Sem().Get(var)->StorageClass(), ast::StorageClass::kFunction); } TEST_F(TypeDeterminerTest, StorageClass_DoesNotSetOnConst) { @@ -1501,7 +1502,7 @@ TEST_F(TypeDeterminerTest, StorageClass_DoesNotSetOnConst) { EXPECT_TRUE(td()->Determine()) << td()->error(); - EXPECT_EQ(var->storage_class(), ast::StorageClass::kNone); + EXPECT_EQ(Sem().Get(var)->StorageClass(), ast::StorageClass::kNone); } TEST_F(TypeDeterminerTest, StorageClass_NonFunctionClassError) { diff --git a/src/validator/validator_impl.cc b/src/validator/validator_impl.cc index 0b0b1631cd..265ad84f27 100644 --- a/src/validator/validator_impl.cc +++ b/src/validator/validator_impl.cc @@ -31,6 +31,7 @@ #include "src/ast/uint_literal.h" #include "src/ast/variable_decl_statement.h" #include "src/semantic/expression.h" +#include "src/semantic/variable.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" #include "src/type/bool_type.h" @@ -336,19 +337,26 @@ bool ValidatorImpl::ValidateConstructedTypes( bool ValidatorImpl::ValidateGlobalVariables( const ast::VariableList& global_vars) { for (auto* var : global_vars) { + auto* sem = program_->Sem().Get(var); + if (!sem) { + add_error(var->source(), "no semantic information for variable '" + + program_->Symbols().NameFor(var->symbol()) + + "'"); + return false; + } + if (variable_stack_.has(var->symbol())) { add_error(var->source(), "v-0011", "redeclared global identifier '" + program_->Symbols().NameFor(var->symbol()) + "'"); return false; } - if (!var->is_const() && var->storage_class() == ast::StorageClass::kNone) { + if (!var->is_const() && sem->StorageClass() == ast::StorageClass::kNone) { add_error(var->source(), "v-0022", "global variables must have a storage class"); return false; } - if (var->is_const() && - !(var->storage_class() == ast::StorageClass::kNone)) { + if (var->is_const() && !(sem->StorageClass() == ast::StorageClass::kNone)) { add_error(var->source(), "v-global01", "global constants shouldn't have a storage class"); return false; diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 62e52e5ce6..6e4eff6c25 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -47,6 +47,7 @@ #include "src/program_builder.h" #include "src/semantic/expression.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/access_control_type.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" @@ -194,7 +195,8 @@ bool GeneratorImpl::Generate(std::ostream& out) { } void GeneratorImpl::register_global(ast::Variable* global) { - global_variables_.set(global->symbol(), global); + auto* sem = builder_.Sem().Get(global); + global_variables_.set(global->symbol(), sem); } std::string GeneratorImpl::generate_name(const std::string& prefix) { @@ -1134,10 +1136,11 @@ bool GeneratorImpl::EmitExpression(std::ostream& pre, return false; } -bool GeneratorImpl::global_is_in_struct(ast::Variable* var) const { - if (var->HasLocationDecoration() || var->HasBuiltinDecoration()) { - return var->storage_class() == ast::StorageClass::kInput || - var->storage_class() == ast::StorageClass::kOutput; +bool GeneratorImpl::global_is_in_struct(const semantic::Variable* var) const { + if (var->Declaration()->HasLocationDecoration() || + var->Declaration()->HasBuiltinDecoration()) { + return var->StorageClass() == ast::StorageClass::kInput || + var->StorageClass() == ast::StorageClass::kOutput; } return false; } @@ -1146,10 +1149,10 @@ bool GeneratorImpl::EmitIdentifier(std::ostream&, std::ostream& out, ast::IdentifierExpression* expr) { auto* ident = expr->As(); - ast::Variable* var = nullptr; + const semantic::Variable* var = nullptr; if (global_variables_.get(ident->symbol(), &var)) { if (global_is_in_struct(var)) { - auto var_type = var->storage_class() == ast::StorageClass::kInput + auto var_type = var->StorageClass() == ast::StorageClass::kInput ? VarType::kIn : VarType::kOut; auto name = current_ep_var_name(var_type); @@ -1230,14 +1233,14 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct( const semantic::Function* func) { for (auto data : func->ReferencedLocationVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kInput) { return true; } } for (auto data : func->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kInput) { return true; } } @@ -1248,14 +1251,14 @@ bool GeneratorImpl::has_referenced_out_var_needing_struct( const semantic::Function* func) { for (auto data : func->ReferencedLocationVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput) { + if (var->StorageClass() == ast::StorageClass::kOutput) { return true; } } for (auto data : func->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput) { + if (var->StorageClass() == ast::StorageClass::kOutput) { return true; } } @@ -1266,16 +1269,16 @@ bool GeneratorImpl::has_referenced_var_needing_struct( const semantic::Function* func) { for (auto data : func->ReferencedLocationVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput || - var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kOutput || + var->StorageClass() == ast::StorageClass::kInput) { return true; } } for (auto data : func->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput || - var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kOutput || + var->StorageClass() == ast::StorageClass::kInput) { return true; } } @@ -1407,51 +1410,54 @@ bool GeneratorImpl::EmitEntryPointData( for (auto data : func_sem->ReferencedLocationVariables()) { auto* var = data.first; + auto* decl = var->Declaration(); auto* deco = data.second; - if (var->storage_class() == ast::StorageClass::kInput) { - in_variables.push_back({var, deco}); - } else if (var->storage_class() == ast::StorageClass::kOutput) { - outvariables.push_back({var, deco}); + if (var->StorageClass() == ast::StorageClass::kInput) { + in_variables.push_back({decl, deco}); + } else if (var->StorageClass() == ast::StorageClass::kOutput) { + outvariables.push_back({decl, deco}); } } for (auto data : func_sem->ReferencedBuiltinVariables()) { auto* var = data.first; + auto* decl = var->Declaration(); auto* deco = data.second; - if (var->storage_class() == ast::StorageClass::kInput) { - in_variables.push_back({var, deco}); - } else if (var->storage_class() == ast::StorageClass::kOutput) { - outvariables.push_back({var, deco}); + if (var->StorageClass() == ast::StorageClass::kInput) { + in_variables.push_back({decl, deco}); + } else if (var->StorageClass() == ast::StorageClass::kOutput) { + outvariables.push_back({decl, deco}); } } bool emitted_uniform = false; for (auto data : func_sem->ReferencedUniformVariables()) { auto* var = data.first; + auto* decl = var->Declaration(); // 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 // set. https://bugs.chromium.org/p/tint/issues/detail?id=104 auto* binding = data.second.binding; if (binding == nullptr) { error_ = "unable to find binding information for uniform: " + - builder_.Symbols().NameFor(var->symbol()); + builder_.Symbols().NameFor(decl->symbol()); return false; } // auto* set = data.second.set; // If the global has already been emitted we skip it, it's been emitted by // a previous entry point. - if (emitted_globals.count(var->symbol()) != 0) { + if (emitted_globals.count(decl->symbol()) != 0) { continue; } - emitted_globals.insert(var->symbol()); + emitted_globals.insert(decl->symbol()); - auto* type = var->type()->UnwrapIfNeeded(); + auto* type = decl->type()->UnwrapIfNeeded(); if (auto* strct = type->As()) { out << "ConstantBuffer<" << builder_.Symbols().NameFor(strct->symbol()) - << "> " << builder_.Symbols().NameFor(var->symbol()) + << "> " << builder_.Symbols().NameFor(decl->symbol()) << " : register(b" << binding->value() << ");" << std::endl; } else { // TODO(dsinclair): There is outstanding spec work to require all uniform @@ -1460,7 +1466,7 @@ bool GeneratorImpl::EmitEntryPointData( // is not a block. // Relevant: https://github.com/gpuweb/gpuweb/issues/1004 // https://github.com/gpuweb/gpuweb/issues/1008 - auto name = "cbuffer_" + builder_.Symbols().NameFor(var->symbol()); + auto name = "cbuffer_" + builder_.Symbols().NameFor(decl->symbol()); out << "cbuffer " << name << " : register(b" << binding->value() << ") {" << std::endl; @@ -1469,7 +1475,7 @@ bool GeneratorImpl::EmitEntryPointData( if (!EmitType(out, type, "")) { return false; } - out << " " << builder_.Symbols().NameFor(var->symbol()) << ";" + out << " " << builder_.Symbols().NameFor(decl->symbol()) << ";" << std::endl; decrement_indent(); out << "};" << std::endl; @@ -1484,16 +1490,17 @@ bool GeneratorImpl::EmitEntryPointData( bool emitted_storagebuffer = false; for (auto data : func_sem->ReferencedStoragebufferVariables()) { auto* var = data.first; + auto* decl = var->Declaration(); auto* binding = data.second.binding; // If the global has already been emitted we skip it, it's been emitted by // a previous entry point. - if (emitted_globals.count(var->symbol()) != 0) { + if (emitted_globals.count(decl->symbol()) != 0) { continue; } - emitted_globals.insert(var->symbol()); + emitted_globals.insert(decl->symbol()); - auto* ac = var->type()->As(); + auto* ac = decl->type()->As(); if (ac == nullptr) { error_ = "access control type required for storage buffer"; return false; @@ -1502,7 +1509,7 @@ bool GeneratorImpl::EmitEntryPointData( if (ac->IsReadWrite()) { out << "RW"; } - out << "ByteAddressBuffer " << builder_.Symbols().NameFor(var->symbol()) + out << "ByteAddressBuffer " << builder_.Symbols().NameFor(decl->symbol()) << " : register(u" << binding->value() << ");" << std::endl; emitted_storagebuffer = true; } @@ -2069,11 +2076,11 @@ bool GeneratorImpl::is_storage_buffer_access( // Check if this is a storage buffer variable if (auto* ident = expr->structure()->As()) { - ast::Variable* var = nullptr; + const semantic::Variable* var = nullptr; if (!global_variables_.get(ident->symbol(), &var)) { return false; } - return var->storage_class() == ast::StorageClass::kStorage; + return var->StorageClass() == ast::StorageClass::kStorage; } else if (auto* member = structure->As()) { return is_storage_buffer_access(member); } else if (auto* array = structure->As()) { diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 24089ae64f..0c175fc86c 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -338,7 +338,7 @@ class GeneratorImpl { /// Checks if the global variable is in an input or output struct /// @param var the variable to check /// @returns true if the global is in an input or output struct - bool global_is_in_struct(ast::Variable* var) const; + bool global_is_in_struct(const semantic::Variable* var) const; /// Creates a text string representing the index into a storage buffer /// @param pre the pre stream /// @param expr the expression to use as the index @@ -400,7 +400,7 @@ class GeneratorImpl { Symbol current_ep_sym_; bool generating_entry_point_ = false; uint32_t loop_emission_counter_ = 0; - ScopeStack global_variables_; + ScopeStack global_variables_; std::unordered_map ep_sym_to_in_data_; std::unordered_map ep_sym_to_out_data_; diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 0f9d249679..be67f4d96c 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -52,6 +52,7 @@ #include "src/program.h" #include "src/semantic/expression.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/access_control_type.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" @@ -120,7 +121,8 @@ bool GeneratorImpl::Generate() { out_ << "#include " << std::endl << std::endl; for (auto* global : program_->AST().GlobalVariables()) { - global_variables_.set(global->symbol(), global); + auto* sem = program_->Sem().Get(global); + global_variables_.set(global->symbol(), sem); } for (auto* const ty : program_->AST().ConstructedTypes()) { @@ -508,14 +510,14 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { auto* func_sem = program_->Sem().Get(func); for (const auto& data : func_sem->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() != ast::StorageClass::kInput) { + if (var->StorageClass() != ast::StorageClass::kInput) { continue; } if (!first) { out_ << ", "; } first = false; - out_ << program_->Symbols().NameFor(var->symbol()); + out_ << program_->Symbols().NameFor(var->Declaration()->symbol()); } for (const auto& data : func_sem->ReferencedUniformVariables()) { @@ -524,7 +526,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { out_ << ", "; } first = false; - out_ << program_->Symbols().NameFor(var->symbol()); + out_ << program_->Symbols().NameFor(var->Declaration()->symbol()); } for (const auto& data : func_sem->ReferencedStoragebufferVariables()) { @@ -533,7 +535,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { out_ << ", "; } first = false; - out_ << program_->Symbols().NameFor(var->symbol()); + out_ << program_->Symbols().NameFor(var->Declaration()->symbol()); } const auto& params = expr->params(); @@ -1033,10 +1035,10 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) { auto* var = data.first; auto* deco = data.second; - if (var->storage_class() == ast::StorageClass::kInput) { - in_locations.push_back({var, deco->value()}); - } else if (var->storage_class() == ast::StorageClass::kOutput) { - out_variables.push_back({var, deco}); + if (var->StorageClass() == ast::StorageClass::kInput) { + in_locations.push_back({var->Declaration(), deco->value()}); + } else if (var->StorageClass() == ast::StorageClass::kOutput) { + out_variables.push_back({var->Declaration(), deco}); } } @@ -1044,8 +1046,8 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) { auto* var = data.first; auto* deco = data.second; - if (var->storage_class() == ast::StorageClass::kOutput) { - out_variables.push_back({var, deco}); + if (var->StorageClass() == ast::StorageClass::kOutput) { + out_variables.push_back({var->Declaration(), deco}); } } @@ -1191,7 +1193,7 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) { auto* func_sem = program_->Sem().Get(func); for (auto data : func_sem->ReferencedLocationVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kInput) { + if (var->StorageClass() == ast::StorageClass::kInput) { return true; } } @@ -1203,14 +1205,14 @@ bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) { for (auto data : func_sem->ReferencedLocationVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput) { + if (var->StorageClass() == ast::StorageClass::kOutput) { return true; } } for (auto data : func_sem->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() == ast::StorageClass::kOutput) { + if (var->StorageClass() == ast::StorageClass::kOutput) { return true; } } @@ -1308,7 +1310,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, for (const auto& data : func_sem->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() != ast::StorageClass::kInput) { + if (var->StorageClass() != ast::StorageClass::kInput) { continue; } if (!first) { @@ -1317,10 +1319,10 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, first = false; out_ << "thread "; - if (!EmitType(var->type(), "")) { + if (!EmitType(var->Declaration()->type(), "")) { return false; } - out_ << "& " << program_->Symbols().NameFor(var->symbol()); + out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol()); } for (const auto& data : func_sem->ReferencedUniformVariables()) { @@ -1332,10 +1334,10 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, out_ << "constant "; // TODO(dsinclair): Can arrays be uniform? If so, fix this ... - if (!EmitType(var->type(), "")) { + if (!EmitType(var->Declaration()->type(), "")) { return false; } - out_ << "& " << program_->Symbols().NameFor(var->symbol()); + out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol()); } for (const auto& data : func_sem->ReferencedStoragebufferVariables()) { @@ -1345,7 +1347,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, } first = false; - auto* ac = var->type()->As(); + auto* ac = var->Declaration()->type()->As(); if (ac == nullptr) { error_ = "invalid type for storage buffer, expected access control"; return false; @@ -1358,7 +1360,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func, if (!EmitType(ac->type(), "")) { return false; } - out_ << "& " << program_->Symbols().NameFor(var->symbol()); + out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol()); } for (auto* v : func->params()) { @@ -1447,7 +1449,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { for (auto data : func_sem->ReferencedBuiltinVariables()) { auto* var = data.first; - if (var->storage_class() != ast::StorageClass::kInput) { + if (var->StorageClass() != ast::StorageClass::kInput) { continue; } @@ -1458,7 +1460,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { auto* builtin = data.second; - if (!EmitType(var->type(), "")) { + if (!EmitType(var->Declaration()->type(), "")) { return false; } @@ -1467,8 +1469,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { error_ = "unknown builtin"; return false; } - out_ << " " << program_->Symbols().NameFor(var->symbol()) << " [[" << attr - << "]]"; + out_ << " " << program_->Symbols().NameFor(var->Declaration()->symbol()) + << " [[" << attr << "]]"; } for (auto data : func_sem->ReferencedUniformVariables()) { @@ -1484,7 +1486,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { auto* binding = data.second.binding; if (binding == nullptr) { error_ = "unable to find binding information for uniform: " + - program_->Symbols().NameFor(var->symbol()); + program_->Symbols().NameFor(var->Declaration()->symbol()); return false; } // auto* set = data.second.set; @@ -1492,11 +1494,11 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { out_ << "constant "; // TODO(dsinclair): Can you have a uniform array? If so, this needs to be // updated to handle arrays property. - if (!EmitType(var->type(), "")) { + if (!EmitType(var->Declaration()->type(), "")) { return false; } - out_ << "& " << program_->Symbols().NameFor(var->symbol()) << " [[buffer(" - << binding->value() << ")]]"; + out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol()) + << " [[buffer(" << binding->value() << ")]]"; } for (auto data : func_sem->ReferencedStoragebufferVariables()) { @@ -1512,7 +1514,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { auto* binding = data.second.binding; // auto* set = data.second.set; - auto* ac = var->type()->As(); + auto* ac = var->Declaration()->type()->As(); if (ac == nullptr) { error_ = "invalid type for storage buffer, expected access control"; return false; @@ -1525,8 +1527,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { if (!EmitType(ac->type(), "")) { return false; } - out_ << "& " << program_->Symbols().NameFor(var->symbol()) << " [[buffer(" - << binding->value() << ")]]"; + out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol()) + << " [[buffer(" << binding->value() << ")]]"; } out_ << ") {" << std::endl; @@ -1563,23 +1565,23 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) { return true; } -bool GeneratorImpl::global_is_in_struct(ast::Variable* var) const { +bool GeneratorImpl::global_is_in_struct(const semantic::Variable* var) const { bool in_or_out_struct_has_location = - var != nullptr && var->HasLocationDecoration() && - (var->storage_class() == ast::StorageClass::kInput || - var->storage_class() == ast::StorageClass::kOutput); + var != nullptr && var->Declaration()->HasLocationDecoration() && + (var->StorageClass() == ast::StorageClass::kInput || + var->StorageClass() == ast::StorageClass::kOutput); bool in_struct_has_builtin = - var != nullptr && var->HasBuiltinDecoration() && - var->storage_class() == ast::StorageClass::kOutput; + var != nullptr && var->Declaration()->HasBuiltinDecoration() && + var->StorageClass() == ast::StorageClass::kOutput; return in_or_out_struct_has_location || in_struct_has_builtin; } bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) { auto* ident = expr->As(); - ast::Variable* var = nullptr; + const semantic::Variable* var = nullptr; if (global_variables_.get(ident->symbol(), &var)) { if (global_is_in_struct(var)) { - auto var_type = var->storage_class() == ast::StorageClass::kInput + auto var_type = var->StorageClass() == ast::StorageClass::kInput ? VarType::kIn : VarType::kOut; auto name = current_ep_var_name(var_type); @@ -1623,7 +1625,7 @@ bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) { // will be turned into assignments. for (auto* s : *(stmt->body())) { if (auto* decl = s->As()) { - if (!EmitVariable(decl->variable(), true)) { + if (!EmitVariable(program_->Sem().Get(decl->variable()), true)) { return false; } } @@ -1839,7 +1841,8 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) { return EmitSwitch(s); } if (auto* v = stmt->As()) { - return EmitVariable(v->variable(), false); + auto* var = program_->Sem().Get(v->variable()); + return EmitVariable(var, false); } error_ = "unknown statement type: " + program_->str(stmt); @@ -2078,35 +2081,38 @@ bool GeneratorImpl::EmitUnaryOp(ast::UnaryOpExpression* expr) { return true; } -bool GeneratorImpl::EmitVariable(ast::Variable* var, bool skip_constructor) { +bool GeneratorImpl::EmitVariable(const semantic::Variable* var, + bool skip_constructor) { make_indent(); + auto* decl = var->Declaration(); + // TODO(dsinclair): Handle variable decorations - if (!var->decorations().empty()) { + if (!decl->decorations().empty()) { error_ = "Variable decorations are not handled yet"; return false; } - if (var->is_const()) { + if (decl->is_const()) { out_ << "const "; } - if (!EmitType(var->type(), program_->Symbols().NameFor(var->symbol()))) { + if (!EmitType(decl->type(), program_->Symbols().NameFor(decl->symbol()))) { return false; } - if (!var->type()->Is()) { - out_ << " " << program_->Symbols().NameFor(var->symbol()); + if (!decl->type()->Is()) { + out_ << " " << program_->Symbols().NameFor(decl->symbol()); } if (!skip_constructor) { out_ << " = "; - if (var->constructor() != nullptr) { - if (!EmitExpression(var->constructor())) { + if (decl->constructor() != nullptr) { + if (!EmitExpression(decl->constructor())) { return false; } - } else if (var->storage_class() == ast::StorageClass::kPrivate || - var->storage_class() == ast::StorageClass::kFunction || - var->storage_class() == ast::StorageClass::kNone || - var->storage_class() == ast::StorageClass::kOutput) { - if (!EmitZeroValue(var->type())) { + } else if (var->StorageClass() == ast::StorageClass::kPrivate || + var->StorageClass() == ast::StorageClass::kFunction || + var->StorageClass() == ast::StorageClass::kNone || + var->StorageClass() == ast::StorageClass::kOutput) { + if (!EmitZeroValue(decl->type())) { return false; } } diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h index d67b54dcda..0aafa6261e 100644 --- a/src/writer/msl/generator_impl.h +++ b/src/writer/msl/generator_impl.h @@ -221,7 +221,7 @@ class GeneratorImpl : public TextGenerator { /// @param var the variable to generate /// @param skip_constructor set true if the constructor should be skipped /// @returns true if the variable was emitted - bool EmitVariable(ast::Variable* var, bool skip_constructor); + bool EmitVariable(const semantic::Variable* var, bool skip_constructor); /// Handles generating a program scope constant variable /// @param var the variable to emit /// @returns true if the variable was emitted @@ -256,7 +256,7 @@ class GeneratorImpl : public TextGenerator { /// Checks if the global variable is in an input or output struct /// @param var the variable to check /// @returns true if the global is in an input or output struct - bool global_is_in_struct(ast::Variable* var) const; + bool global_is_in_struct(const semantic::Variable* var) const; /// Converts a builtin to an attribute name /// @param builtin the builtin to convert @@ -283,7 +283,7 @@ class GeneratorImpl : public TextGenerator { } Namer namer_; - ScopeStack global_variables_; + ScopeStack global_variables_; Symbol current_ep_sym_; bool generating_entry_point_ = false; const Program* program_ = nullptr; diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 645c6b4e1b..79c71e864b 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -61,6 +61,7 @@ #include "src/program.h" #include "src/semantic/expression.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/access_control_type.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" @@ -462,15 +463,15 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) { for (const auto* var : func_sem->ReferencedModuleVariables()) { // 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->storage_class() != ast::StorageClass::kInput && - var->storage_class() != ast::StorageClass::kOutput) { + if (var->StorageClass() != ast::StorageClass::kInput && + var->StorageClass() != ast::StorageClass::kOutput) { continue; } uint32_t var_id; - if (!scope_stack_.get(var->symbol(), &var_id)) { + if (!scope_stack_.get(var->Declaration()->symbol(), &var_id)) { error_ = "unable to find ID for global variable: " + - builder_.Symbols().NameFor(var->symbol()); + builder_.Symbols().NameFor(var->Declaration()->symbol()); return false; } @@ -700,6 +701,8 @@ bool Builder::GenerateStore(uint32_t to, uint32_t from) { } bool Builder::GenerateGlobalVariable(ast::Variable* var) { + auto* sem = builder_.Sem().Get(var); + uint32_t init_id = 0; if (var->has_constructor()) { if (!var->constructor()->Is()) { @@ -731,9 +734,9 @@ bool Builder::GenerateGlobalVariable(ast::Variable* var) { auto result = result_op(); auto var_id = result.to_i(); - auto sc = var->storage_class() == ast::StorageClass::kNone + auto sc = sem->StorageClass() == ast::StorageClass::kNone ? ast::StorageClass::kPrivate - : var->storage_class(); + : sem->StorageClass(); type::Pointer pt(var->type(), sc); auto type_id = GenerateTypeIfNeeded(&pt); @@ -799,9 +802,9 @@ bool Builder::GenerateGlobalVariable(ast::Variable* var) { return 0; } ops.push_back(Operand::Int(init_id)); - } else if (var->storage_class() == ast::StorageClass::kPrivate || - var->storage_class() == ast::StorageClass::kNone || - var->storage_class() == ast::StorageClass::kOutput) { + } else if (sem->StorageClass() == ast::StorageClass::kPrivate || + sem->StorageClass() == ast::StorageClass::kNone || + sem->StorageClass() == ast::StorageClass::kOutput) { ast::NullLiteral nl(Source{}, type); init_id = GenerateLiteralIfNeeded(var, &nl); if (init_id == 0) { diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc index 7410675778..3025ffddc9 100644 --- a/src/writer/wgsl/generator_impl.cc +++ b/src/writer/wgsl/generator_impl.cc @@ -59,6 +59,7 @@ #include "src/ast/workgroup_decoration.h" #include "src/program.h" #include "src/semantic/function.h" +#include "src/semantic/variable.h" #include "src/type/access_control_type.h" #include "src/type/alias_type.h" #include "src/type/array_type.h" @@ -148,7 +149,7 @@ bool GeneratorImpl::GenerateEntryPoint(ast::PipelineStage stage, bool found_func_variable = false; for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) { - if (!EmitVariable(var)) { + if (!EmitVariable(var->Declaration())) { return false; } found_func_variable = true; @@ -583,9 +584,11 @@ bool GeneratorImpl::EmitStructType(const type::Struct* str) { } bool GeneratorImpl::EmitVariable(ast::Variable* var) { + auto* sem = program_->Sem().Get(var); + make_indent(); - if (!var->decorations().empty() && !EmitVariableDecorations(var)) { + if (!var->decorations().empty() && !EmitVariableDecorations(sem)) { return false; } @@ -593,9 +596,9 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) { out_ << "const"; } else { out_ << "var"; - if (var->storage_class() != ast::StorageClass::kNone && - var->storage_class() != ast::StorageClass::kFunction) { - out_ << "<" << var->storage_class() << ">"; + if (sem->StorageClass() != ast::StorageClass::kNone && + sem->StorageClass() != ast::StorageClass::kFunction) { + out_ << "<" << sem->StorageClass() << ">"; } } @@ -615,10 +618,12 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) { return true; } -bool GeneratorImpl::EmitVariableDecorations(ast::Variable* var) { +bool GeneratorImpl::EmitVariableDecorations(const semantic::Variable* var) { + auto* decl = var->Declaration(); + out_ << "[["; bool first = true; - for (auto* deco : var->decorations()) { + for (auto* deco : decl->decorations()) { if (!first) { out_ << ", "; } diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h index fa4603b3f2..34ad398261 100644 --- a/src/writer/wgsl/generator_impl.h +++ b/src/writer/wgsl/generator_impl.h @@ -202,7 +202,7 @@ class GeneratorImpl : public TextGenerator { /// Handles generating variable decorations /// @param var the decorated variable /// @returns true if the variable decoration was emitted - bool EmitVariableDecorations(ast::Variable* var); + bool EmitVariableDecorations(const semantic::Variable* var); private: Program const* const program_;