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>
This commit is contained in:
parent
401b96b9bb
commit
b17aea159c
1
BUILD.gn
1
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",
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -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<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 @@ 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;
|
||||
}
|
||||
|
|
|
@ -84,7 +84,7 @@ class Variable : public Castable<Variable, Node> {
|
|||
/// 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<Variable, Node> {
|
|||
/// @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<Variable, Node> {
|
|||
|
||||
/// @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<Variable, Node> {
|
|||
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
|
||||
|
|
|
@ -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);
|
||||
|
|
|
@ -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<EntryPoint> 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<EntryPoint> 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<ResourceBinding> 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<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 @@ std::vector<ResourceBinding> 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<ResourceBinding> 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<ResourceBinding> 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<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
|
|||
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 @@ std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
|
|||
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 @@ std::vector<ResourceBinding> 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<type::Texture>();
|
||||
auto* texture_type = decl->type()->UnwrapIfNeeded()->As<type::Texture>();
|
||||
switch (texture_type->dim()) {
|
||||
case type::TextureDimension::k1d:
|
||||
entry.dim = ResourceBinding::TextureDimension::k1d;
|
||||
|
|
|
@ -38,7 +38,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithoutConstructor) {
|
|||
|
||||
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 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithConstructor) {
|
|||
|
||||
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 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration) {
|
|||
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 @@ 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<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);
|
||||
|
|
|
@ -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<Function, Node> {
|
||||
public:
|
||||
|
@ -54,8 +55,8 @@ class Function : public Castable<Function, Node> {
|
|||
/// @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 @@ class Function : public Castable<Function, Node> {
|
|||
/// 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 @@ class Function : public Castable<Function, Node> {
|
|||
}
|
||||
/// 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 @@ class Function : public Castable<Function, Node> {
|
|||
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_;
|
||||
};
|
||||
|
||||
|
|
|
@ -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<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(std::vector<ast::Variable*> referenced_module_vars,
|
|||
|
||||
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 @@ Function::ReferencedLocationVariables() const {
|
|||
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 @@ Function::ReferencedUniformVariables() const {
|
|||
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 @@ Function::ReferencedStoragebufferVariables() const {
|
|||
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 @@ Function::ReferencedBuiltinVariables() const {
|
|||
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 @@ bool Function::HasAncestorEntryPoint(Symbol symbol) const {
|
|||
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 @@ 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<ast::BindingDecoration>()) {
|
||||
binding = b;
|
||||
}
|
||||
|
@ -196,12 +197,12 @@ Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
|
|||
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 @@ 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<ast::BindingDecoration>()) {
|
||||
binding = b;
|
||||
} else if (auto* s = deco->As<ast::GroupDecoration>()) {
|
||||
|
|
|
@ -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
|
|
@ -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
|
||||
};
|
||||
|
||||
|
|
|
@ -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_
|
|
@ -58,7 +58,6 @@ Transform::Output EmitVertexPointSize::Run(const Program* in) {
|
|||
false, // is_const
|
||||
nullptr, // constructor
|
||||
ast::VariableDecorationList{
|
||||
// decorations
|
||||
out.create<ast::BuiltinDecoration>(Source{},
|
||||
ast::Builtin::kPointSize),
|
||||
});
|
||||
|
|
|
@ -70,7 +70,7 @@ ast::Variable* clone_variable_with_new_name(CloneContext* ctx,
|
|||
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::Variable* FirstIndexOffset::AddUniformBuffer(ProgramBuilder* dst) {
|
|||
ast::VariableDecorationList{
|
||||
dst->create<ast::BindingDecoration>(Source{}, binding_),
|
||||
dst->create<ast::GroupDecoration>(Source{}, group_),
|
||||
}); // decorations
|
||||
});
|
||||
|
||||
dst->AST().AddGlobalVariable(idx_var);
|
||||
|
||||
|
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
|
@ -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<type::StorageTexture*> storage_textures;
|
||||
|
||||
for (auto& it : builder_->Types().types()) {
|
||||
if (auto* storage =
|
||||
it.second->UnwrapIfNeeded()->As<type::StorageTexture>()) {
|
||||
|
@ -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<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 @@ 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<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::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<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 @@ 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<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) {}
|
||||
|
||||
|
|
|
@ -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<T> &() const { return vector; }
|
||||
|
@ -108,6 +92,16 @@ class TypeDeterminer {
|
|||
std::unordered_set<T> set;
|
||||
};
|
||||
|
||||
/// Structure holding semantic information about a function.
|
||||
/// Used to build the semantic::Function nodes at the end of resolving.
|
||||
struct 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<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 @@ 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<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.
|
||||
|
|
|
@ -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) {
|
||||
|
|
|
@ -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;
|
||||
|
|
|
@ -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::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 @@ 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<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 @@ 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<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 @@ 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::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>()) {
|
||||
|
|
|
@ -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<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_;
|
||||
|
||||
|
|
|
@ -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 <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 @@ 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<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 @@ 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<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 @@ 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::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 @@ bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
|
|||
// 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 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
|
|||
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 @@ 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<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;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -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<ast::Variable*> global_variables_;
|
||||
ScopeStack<const semantic::Variable*> global_variables_;
|
||||
Symbol current_ep_sym_;
|
||||
bool generating_entry_point_ = false;
|
||||
const Program* program_ = nullptr;
|
||||
|
|
|
@ -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<ast::ConstructorExpression>()) {
|
||||
|
@ -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) {
|
||||
|
|
|
@ -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_ << ", ";
|
||||
}
|
||||
|
|
|
@ -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_;
|
||||
|
|
Loading…
Reference in New Issue