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 <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/BUILD.gn b/BUILD.gn
index 8e15afc..12578bb 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -385,6 +385,7 @@
"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 56a1efb..f79ae1a 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -199,6 +199,7 @@
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 cc0e676..40056a4 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 @@
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 @@
}
Variable* Variable::Clone(CloneContext* ctx) const {
- return ctx->dst->create<Variable>(ctx->Clone(source()), ctx->Clone(symbol_),
- storage_class(), ctx->Clone(type()),
- is_const_, ctx->Clone(constructor()),
- ctx->Clone(decorations_));
+ return ctx->dst->create<Variable>(
+ 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 @@
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 fb337b1..2d18b14 100644
--- a/src/ast/variable.h
+++ b/src/ast/variable.h
@@ -84,7 +84,7 @@
/// 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 @@
/// @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 @@
/// @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 @@
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 f700381..9a9f6be 100644
--- a/src/ast/variable_test.cc
+++ b/src/ast/variable_test.cc
@@ -30,7 +30,7 @@
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 @@
"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 @@
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 ccfe471..ca4bb97 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 @@
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 @@
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 @@
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<type::AccessControl>()) {
+ auto* var = ruv.first;
+ auto* decl = var->Declaration();
+ auto binding_info = ruv.second;
+
+ if (!decl->type()->Is<type::AccessControl>()) {
continue;
}
- auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+ auto* unwrapped_type = decl->type()->UnwrapIfNeeded();
auto* str = unwrapped_type->As<type::Struct>();
if (str == nullptr) {
@@ -206,10 +209,11 @@
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 @@
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 @@
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 @@
auto* func_sem = program_->Sem().Get(func);
std::vector<ResourceBinding> 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<type::AccessControl>();
+ auto* ac_type = decl->type()->As<type::AccessControl>();
if (ac_type == nullptr) {
continue;
}
@@ -328,14 +327,15 @@
continue;
}
- if (!var->type()->UnwrapIfNeeded()->Is<type::Struct>()) {
+ if (!decl->type()->UnwrapIfNeeded()->Is<type::Struct>()) {
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 @@
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<type::Texture>();
+ auto* texture_type = decl->type()->UnwrapIfNeeded()->As<type::Texture>();
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 9c83e42..625abb5 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 @@
EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
EXPECT_TRUE(e->type()->Is<type::F32>());
- 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 @@
EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
EXPECT_TRUE(e->type()->Is<type::F32>());
- 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 @@
EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
ASSERT_NE(e->type(), nullptr);
EXPECT_TRUE(e->type()->Is<type::F32>());
- 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 @@
EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
ASSERT_NE(e->type(), nullptr);
EXPECT_TRUE(e->type()->Is<type::F32>());
- 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 0c75113..70708d9 100644
--- a/src/semantic/function.h
+++ b/src/semantic/function.h
@@ -28,7 +28,6 @@
namespace ast {
class BindingDecoration;
class GroupDecoration;
-class Variable;
class LocationDecoration;
class BuiltinDecoration;
} // namespace ast
@@ -38,6 +37,8 @@
namespace semantic {
+class Variable;
+
/// Function holds the semantic information for function nodes.
class Function : public Castable<Function, Node> {
public:
@@ -54,8 +55,8 @@
/// @param local_referenced_module_vars the locally referenced module
/// variables
/// @param ancestor_entry_points the ancestor entry points
- explicit Function(std::vector<ast::Variable*> referenced_module_vars,
- std::vector<ast::Variable*> local_referenced_module_vars,
+ explicit Function(std::vector<const Variable*> referenced_module_vars,
+ std::vector<const Variable*> local_referenced_module_vars,
std::vector<Symbol> ancestor_entry_points);
/// Destructor
@@ -64,11 +65,11 @@
/// Note: If this function calls other functions, the return will also include
/// all of the referenced variables from the callees.
/// @returns the referenced module variables
- const std::vector<ast::Variable*>& ReferencedModuleVariables() const {
+ const std::vector<const Variable*>& ReferencedModuleVariables() const {
return referenced_module_vars_;
}
/// @returns the locally referenced module variables
- const std::vector<ast::Variable*>& LocalReferencedModuleVariables() const {
+ const std::vector<const Variable*>& LocalReferencedModuleVariables() const {
return local_referenced_module_vars_;
}
/// @returns the ancestor entry points
@@ -77,53 +78,53 @@
}
/// Retrieves any referenced location variables
/// @returns the <variable, decoration> pair.
- const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
+ const std::vector<std::pair<const Variable*, ast::LocationDecoration*>>
ReferencedLocationVariables() const;
/// Retrieves any referenced builtin variables
/// @returns the <variable, decoration> pair.
- const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+ const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
ReferencedBuiltinVariables() const;
/// Retrieves any referenced uniform variables. Note, the variables must be
/// decorated with both binding and group decorations.
/// @returns the referenced uniforms
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedUniformVariables() const;
/// Retrieves any referenced storagebuffer variables. Note, the variables
/// must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedStoragebufferVariables() const;
/// Retrieves any referenced regular Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedSamplerVariables() const;
/// Retrieves any referenced comparison Sampler variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced storagebuffers
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedComparisonSamplerVariables() const;
/// Retrieves any referenced sampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedSampledTextureVariables() const;
/// Retrieves any referenced multisampled textures variables. Note, the
/// variables must be decorated with both binding and group decorations.
/// @returns the referenced sampled textures
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedMultisampledTextureVariables() const;
/// Retrieves any locally referenced builtin variables
/// @returns the <variable, decoration> pairs.
- const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+ const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
LocalReferencedBuiltinVariables() const;
/// Checks if the given entry point is an ancestor
@@ -132,13 +133,13 @@
bool HasAncestorEntryPoint(Symbol sym) const;
private:
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedSamplerVariablesImpl(type::SamplerKind kind) const;
- const std::vector<std::pair<ast::Variable*, BindingInfo>>
+ const std::vector<std::pair<const Variable*, BindingInfo>>
ReferencedSampledTextureVariablesImpl(bool multisampled) const;
- std::vector<ast::Variable*> const referenced_module_vars_;
- std::vector<ast::Variable*> const local_referenced_module_vars_;
+ std::vector<const Variable*> const referenced_module_vars_;
+ std::vector<const Variable*> const local_referenced_module_vars_;
std::vector<Symbol> const ancestor_entry_points_;
};
diff --git a/src/semantic/sem_function.cc b/src/semantic/sem_function.cc
index 5613697..26768ec 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 @@
namespace tint {
namespace semantic {
-Function::Function(std::vector<ast::Variable*> referenced_module_vars,
- std::vector<ast::Variable*> local_referenced_module_vars,
+Function::Function(std::vector<const Variable*> referenced_module_vars,
+ std::vector<const Variable*> local_referenced_module_vars,
std::vector<Symbol> ancestor_entry_points)
: referenced_module_vars_(std::move(referenced_module_vars)),
local_referenced_module_vars_(std::move(local_referenced_module_vars)),
@@ -38,12 +39,12 @@
Function::~Function() = default;
-const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
+const std::vector<std::pair<const Variable*, ast::LocationDecoration*>>
Function::ReferencedLocationVariables() const {
- std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>> ret;
+ std::vector<std::pair<const Variable*, ast::LocationDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
- for (auto* deco : var->decorations()) {
+ for (auto* deco : var->Declaration()->decorations()) {
if (auto* location = deco->As<ast::LocationDecoration>()) {
ret.push_back({var, location});
break;
@@ -53,18 +54,18 @@
return ret;
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedUniformVariables() const {
- std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+ std::vector<std::pair<const Variable*, Function::BindingInfo>> 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<ast::BindingDecoration>()) {
binding = b;
} else if (auto* g = deco->As<ast::GroupDecoration>()) {
@@ -80,18 +81,18 @@
return ret;
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedStoragebufferVariables() const {
- std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+ std::vector<std::pair<const Variable*, Function::BindingInfo>> 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<ast::BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<ast::GroupDecoration>()) {
@@ -107,12 +108,12 @@
return ret;
}
-const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
Function::ReferencedBuiltinVariables() const {
- std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
+ std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>> ret;
for (auto* var : ReferencedModuleVariables()) {
- for (auto* deco : var->decorations()) {
+ for (auto* deco : var->Declaration()->decorations()) {
if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
@@ -122,32 +123,32 @@
return ret;
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedSamplerVariables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kSampler);
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedComparisonSamplerVariables() const {
return ReferencedSamplerVariablesImpl(type::SamplerKind::kComparisonSampler);
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedSampledTextureVariables() const {
return ReferencedSampledTextureVariablesImpl(false);
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedMultisampledTextureVariables() const {
return ReferencedSampledTextureVariablesImpl(true);
}
-const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
Function::LocalReferencedBuiltinVariables() const {
- std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
+ std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>> ret;
for (auto* var : LocalReferencedModuleVariables()) {
- for (auto* deco : var->decorations()) {
+ for (auto* deco : var->Declaration()->decorations()) {
if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
ret.push_back({var, builtin});
break;
@@ -166,12 +167,12 @@
return false;
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
- std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+ std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
- auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+ auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded();
auto* sampler = unwrapped_type->As<type::Sampler>();
if (sampler == nullptr || sampler->kind() != kind) {
continue;
@@ -179,7 +180,7 @@
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
- for (auto* deco : var->decorations()) {
+ for (auto* deco : var->Declaration()->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
}
@@ -196,12 +197,12 @@
return ret;
}
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const {
- std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+ std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
for (auto* var : ReferencedModuleVariables()) {
- auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+ auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded();
auto* texture = unwrapped_type->As<type::Texture>();
if (texture == nullptr) {
continue;
@@ -216,7 +217,7 @@
ast::BindingDecoration* binding = nullptr;
ast::GroupDecoration* group = nullptr;
- for (auto* deco : var->decorations()) {
+ for (auto* deco : var->Declaration()->decorations()) {
if (auto* b = deco->As<ast::BindingDecoration>()) {
binding = b;
} else if (auto* s = deco->As<ast::GroupDecoration>()) {
diff --git a/src/semantic/sem_variable.cc b/src/semantic/sem_variable.cc
new file mode 100644
index 0000000..33e84a1
--- /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 ac099a8..58b91c5 100644
--- a/src/semantic/type_mappings.h
+++ b/src/semantic/type_mappings.h
@@ -24,6 +24,7 @@
class Expression;
class Function;
+class Variable;
} // namespace ast
@@ -31,6 +32,7 @@
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 @@
//! @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 0000000..6f15dc7
--- /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 <utility>
+#include <vector>
+
+#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<Variable, Node> {
+ 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 319d144..748021c 100644
--- a/src/transform/emit_vertex_point_size.cc
+++ b/src/transform/emit_vertex_point_size.cc
@@ -58,7 +58,6 @@
false, // is_const
nullptr, // constructor
ast::VariableDecorationList{
- // decorations
out.create<ast::BuiltinDecoration>(Source{},
ast::Builtin::kPointSize),
});
diff --git a/src/transform/first_index_offset.cc b/src/transform/first_index_offset.cc
index fec1068..4ed52ab 100644
--- a/src/transform/first_index_offset.cc
+++ b/src/transform/first_index_offset.cc
@@ -70,7 +70,7 @@
return ctx->dst->create<ast::Variable>(
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::VariableDecorationList{
dst->create<ast::BindingDecoration>(Source{}, binding_),
dst->create<ast::GroupDecoration>(Source{}, group_),
- }); // decorations
+ });
dst->AST().AddGlobalVariable(idx_var);
diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc
index 39585be..f052bd4 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 @@
// 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 @@
// 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::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 6172384..e75bdc8 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 @@
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 @@
}
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<type::StorageTexture*> storage_textures;
+
for (auto& it : builder_->Types().types()) {
if (auto* storage =
it.second->UnwrapIfNeeded()->As<type::StorageTexture>()) {
@@ -122,7 +134,7 @@
}
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 @@
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 @@
}
}
- CreateSemanticFunctions();
-
return true;
}
@@ -170,15 +180,13 @@
}
bool TypeDeterminer::DetermineFunction(ast::Function* func) {
- auto* info = function_infos_.Create<FunctionInfo>(func);
- symbol_to_function_[func->symbol()] = info;
- function_to_info_.emplace(func, info);
-
- current_function_ = info;
+ current_function_ = function_infos_.Create<FunctionInfo>(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 @@
}
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 @@
return true;
}
if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
- 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::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<type::Pointer>()) {
- SetType(expr, var->type());
+ if (var->declaration->is_const()) {
+ SetType(expr, var->declaration->type());
+ } else if (var->declaration->type()->Is<type::Pointer>()) {
+ SetType(expr, var->declaration->type());
} else {
- SetType(expr, builder_->create<type::Pointer>(var->type(),
- var->storage_class()));
+ SetType(expr, builder_->create<type::Pointer>(var->declaration->type(),
+ var->storage_class));
}
set_referenced_from_function_if_needed(var, true);
@@ -1146,6 +1159,13 @@
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 @@
builder_->create<semantic::Expression>(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<semantic::Variable>(var, info->storage_class));
+ }
+
+ auto remap_vars = [&sem](const std::vector<VariableInfo*>& in) {
+ std::vector<const semantic::Variable*> 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<semantic::Function>(
- info->referenced_module_vars,
- info->local_referenced_module_vars,
- info->ancestor_entry_points));
+ sem.Add(func, builder_->create<semantic::Function>(
+ 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 c943a53..1402dda 100644
--- a/src/type_determiner.h
+++ b/src/type_determiner.h
@@ -66,23 +66,6 @@
/// @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 @@
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<T> &() const { return vector; }
@@ -110,16 +94,31 @@
/// 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 {
explicit FunctionInfo(ast::Function* decl);
~FunctionInfo();
ast::Function* const declaration;
- UniqueVector<ast::Variable*> referenced_module_vars;
- UniqueVector<ast::Variable*> local_referenced_module_vars;
+ UniqueVector<VariableInfo*> referenced_module_vars;
+ UniqueVector<VariableInfo*> local_referenced_module_vars;
UniqueVector<Symbol> 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 @@
/// @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 @@
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 @@
ProgramBuilder* builder_;
std::string error_;
- ScopeStack<ast::Variable*> variable_stack_;
+ ScopeStack<VariableInfo*> variable_stack_;
std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_;
std::unordered_map<ast::Function*, FunctionInfo*> function_to_info_;
+ std::unordered_map<ast::Variable*, VariableInfo*> variable_to_info_;
FunctionInfo* current_function_ = nullptr;
+ BlockAllocator<VariableInfo> variable_infos_;
BlockAllocator<FunctionInfo> function_infos_;
// Map from caller functions to callee functions.
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc
index b97f009..96def37 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 @@
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 @@
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 @@
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 @@
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 0b0b163..265ad84 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::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 62e52e5..6e4eff6 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 @@
}
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 @@
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 @@
std::ostream& out,
ast::IdentifierExpression* expr) {
auto* ident = expr->As<ast::IdentifierExpression>();
- 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 @@
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 @@
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 @@
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 @@
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<type::Struct>()) {
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 @@
// 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 @@
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 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<type::AccessControl>();
+ auto* ac = decl->type()->As<type::AccessControl>();
if (ac == nullptr) {
error_ = "access control type required for storage buffer";
return false;
@@ -1502,7 +1509,7 @@
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 @@
// Check if this is a storage buffer variable
if (auto* ident = expr->structure()->As<ast::IdentifierExpression>()) {
- 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<ast::MemberAccessorExpression>()) {
return is_storage_buffer_access(member);
} else if (auto* array = structure->As<ast::ArrayAccessorExpression>()) {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 24089ae..0c175fc 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -338,7 +338,7 @@
/// 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 @@
Symbol current_ep_sym_;
bool generating_entry_point_ = false;
uint32_t loop_emission_counter_ = 0;
- ScopeStack<ast::Variable*> global_variables_;
+ ScopeStack<const semantic::Variable*> global_variables_;
std::unordered_map<Symbol, EntryPointData> ep_sym_to_in_data_;
std::unordered_map<Symbol, EntryPointData> ep_sym_to_out_data_;
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 0f9d249..be67f4d 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 @@
out_ << "#include <metal_stdlib>" << 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
}
first = false;
- auto* ac = var->type()->As<type::AccessControl>();
+ auto* ac = var->Declaration()->type()->As<type::AccessControl>();
if (ac == nullptr) {
error_ = "invalid type for storage buffer, expected access control";
return false;
@@ -1358,7 +1360,7 @@
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 @@
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 @@
auto* builtin = data.second;
- if (!EmitType(var->type(), "")) {
+ if (!EmitType(var->Declaration()->type(), "")) {
return false;
}
@@ -1467,8 +1469,8 @@
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 @@
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 @@
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 @@
auto* binding = data.second.binding;
// auto* set = data.second.set;
- auto* ac = var->type()->As<type::AccessControl>();
+ auto* ac = var->Declaration()->type()->As<type::AccessControl>();
if (ac == nullptr) {
error_ = "invalid type for storage buffer, expected access control";
return false;
@@ -1525,8 +1527,8 @@
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 @@
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::IdentifierExpression>();
- 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 @@
// will be turned into assignments.
for (auto* s : *(stmt->body())) {
if (auto* decl = s->As<ast::VariableDeclStatement>()) {
- if (!EmitVariable(decl->variable(), true)) {
+ if (!EmitVariable(program_->Sem().Get(decl->variable()), true)) {
return false;
}
}
@@ -1839,7 +1841,8 @@
return EmitSwitch(s);
}
if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
- 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 @@
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<type::Array>()) {
- out_ << " " << program_->Symbols().NameFor(var->symbol());
+ if (!decl->type()->Is<type::Array>()) {
+ 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 d67b54d..0aafa62 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -221,7 +221,7 @@
/// @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 @@
/// 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 @@
}
Namer namer_;
- ScopeStack<ast::Variable*> global_variables_;
+ ScopeStack<const semantic::Variable*> 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 645c6b4..79c71e8 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 @@
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::GenerateGlobalVariable(ast::Variable* var) {
+ auto* sem = builder_.Sem().Get(var);
+
uint32_t init_id = 0;
if (var->has_constructor()) {
if (!var->constructor()->Is<ast::ConstructorExpression>()) {
@@ -731,9 +734,9 @@
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 @@
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 7410675..3025ffd 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 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::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 @@
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 @@
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 fa4603b..34ad398 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -202,7 +202,7 @@
/// 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_;
