Google Git
Sign in
dawn / dawn / f3f813eb0c1fe6a3dd87fd048a1f352dee448a05 / . / src / tint / resolver / resolver.cc
blob: 09c4518da1b7bfe7de3c8a6b3889c7774d32576b [file] [log] [blame]
// Copyright 2020 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/tint/resolver/resolver.h"
#include <algorithm>
#include <cmath>
#include <iomanip>
#include <limits>
#include <utility>
#include "src/tint/ast/alias.h"
#include "src/tint/ast/array.h"
#include "src/tint/ast/assignment_statement.h"
#include "src/tint/ast/attribute.h"
#include "src/tint/ast/bitcast_expression.h"
#include "src/tint/ast/break_statement.h"
#include "src/tint/ast/call_statement.h"
#include "src/tint/ast/continue_statement.h"
#include "src/tint/ast/depth_texture.h"
#include "src/tint/ast/disable_validation_attribute.h"
#include "src/tint/ast/discard_statement.h"
#include "src/tint/ast/for_loop_statement.h"
#include "src/tint/ast/id_attribute.h"
#include "src/tint/ast/if_statement.h"
#include "src/tint/ast/internal_attribute.h"
#include "src/tint/ast/interpolate_attribute.h"
#include "src/tint/ast/loop_statement.h"
#include "src/tint/ast/matrix.h"
#include "src/tint/ast/pointer.h"
#include "src/tint/ast/return_statement.h"
#include "src/tint/ast/sampled_texture.h"
#include "src/tint/ast/sampler.h"
#include "src/tint/ast/storage_texture.h"
#include "src/tint/ast/switch_statement.h"
#include "src/tint/ast/traverse_expressions.h"
#include "src/tint/ast/type_name.h"
#include "src/tint/ast/unary_op_expression.h"
#include "src/tint/ast/variable_decl_statement.h"
#include "src/tint/ast/vector.h"
#include "src/tint/ast/while_statement.h"
#include "src/tint/ast/workgroup_attribute.h"
#include "src/tint/resolver/uniformity.h"
#include "src/tint/sem/break_if_statement.h"
#include "src/tint/sem/call.h"
#include "src/tint/sem/for_loop_statement.h"
#include "src/tint/sem/function.h"
#include "src/tint/sem/if_statement.h"
#include "src/tint/sem/index_accessor_expression.h"
#include "src/tint/sem/load.h"
#include "src/tint/sem/loop_statement.h"
#include "src/tint/sem/materialize.h"
#include "src/tint/sem/member_accessor_expression.h"
#include "src/tint/sem/module.h"
#include "src/tint/sem/statement.h"
#include "src/tint/sem/struct.h"
#include "src/tint/sem/switch_statement.h"
#include "src/tint/sem/type_conversion.h"
#include "src/tint/sem/type_initializer.h"
#include "src/tint/sem/variable.h"
#include "src/tint/sem/while_statement.h"
#include "src/tint/type/abstract_float.h"
#include "src/tint/type/abstract_int.h"
#include "src/tint/type/array.h"
#include "src/tint/type/atomic.h"
#include "src/tint/type/depth_multisampled_texture.h"
#include "src/tint/type/depth_texture.h"
#include "src/tint/type/multisampled_texture.h"
#include "src/tint/type/pointer.h"
#include "src/tint/type/reference.h"
#include "src/tint/type/sampled_texture.h"
#include "src/tint/type/sampler.h"
#include "src/tint/type/short_name.h"
#include "src/tint/type/storage_texture.h"
#include "src/tint/utils/defer.h"
#include "src/tint/utils/math.h"
#include "src/tint/utils/reverse.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/utils/string.h"
#include "src/tint/utils/transform.h"
#include "src/tint/utils/vector.h"
namespace tint::resolver {
namespace {
constexpr int64_t kMaxArrayElementCount = 65536;
constexpr uint32_t kMaxStatementDepth = 127;
} // namespace
Resolver::Resolver(ProgramBuilder* builder)
: builder_(builder),
diagnostics_(builder->Diagnostics()),
const_eval_(*builder),
intrinsic_table_(IntrinsicTable::Create(*builder)),
sem_(builder, dependencies_),
validator_(builder,
sem_,
enabled_extensions_,
atomic_composite_info_,
valid_type_storage_layouts_) {}
Resolver::~Resolver() = default;
bool Resolver::Resolve() {
if (builder_->Diagnostics().contains_errors()) {
return false;
}
builder_->Sem().Reserve(builder_->LastAllocatedNodeID());
// Pre-allocate the marked bitset with the total number of AST nodes.
marked_.Resize(builder_->ASTNodes().Count());
if (!DependencyGraph::Build(builder_->AST(), builder_->Symbols(), builder_->Diagnostics(),
dependencies_)) {
return false;
}
bool result = ResolveInternal();
if (!result && !diagnostics_.contains_errors()) {
TINT_ICE(Resolver, diagnostics_) << "resolving failed, but no error was raised";
return false;
}
// Create the semantic module
builder_->Sem().SetModule(builder_->create<sem::Module>(
std::move(dependencies_.ordered_globals), std::move(enabled_extensions_)));
return result;
}
bool Resolver::ResolveInternal() {
Mark(&builder_->AST());
// Process all module-scope declarations in dependency order.
for (auto* decl : dependencies_.ordered_globals) {
Mark(decl);
if (!Switch<bool>(
decl, //
[&](const ast::Enable* e) { return Enable(e); },
[&](const ast::TypeDecl* td) { return TypeDecl(td); },
[&](const ast::Function* func) { return Function(func); },
[&](const ast::Variable* var) { return GlobalVariable(var); },
[&](const ast::StaticAssert* sa) { return StaticAssert(sa); },
[&](Default) {
TINT_UNREACHABLE(Resolver, diagnostics_)
<< "unhandled global declaration: " << decl->TypeInfo().name;
return false;
})) {
return false;
}
}
if (!AllocateOverridableConstantIds()) {
return false;
}
SetShadows();
if (!validator_.PipelineStages(entry_points_)) {
return false;
}
if (!validator_.PushConstants(entry_points_)) {
return false;
}
if (!enabled_extensions_.Contains(ast::Extension::kChromiumDisableUniformityAnalysis)) {
if (!AnalyzeUniformity(builder_, dependencies_)) {
if (kUniformityFailuresAsError) {
return false;
}
}
}
bool result = true;
for (auto* node : builder_->ASTNodes().Objects()) {
if (!marked_[node->node_id.value]) {
TINT_ICE(Resolver, diagnostics_)
<< "AST node '" << node->TypeInfo().name << "' was not reached by the resolver\n"
<< "At: " << node->source << "\n"
<< "Pointer: " << node;
result = false;
}
}
return result;
}
type::Type* Resolver::Type(const ast::Type* ty) {
Mark(ty);
auto* s = Switch(
ty, //
[&](const ast::Void*) { return builder_->create<type::Void>(); },
[&](const ast::Bool*) { return builder_->create<type::Bool>(); },
[&](const ast::I32*) { return builder_->create<type::I32>(); },
[&](const ast::U32*) { return builder_->create<type::U32>(); },
[&](const ast::F16* t) -> type::F16* {
return validator_.CheckF16Enabled(t->source) ? builder_->create<type::F16>() : nullptr;
},
[&](const ast::F32*) { return builder_->create<type::F32>(); },
[&](const ast::Vector* t) -> type::Vector* {
if (!t->type) {
AddError("missing vector element type", t->source.End());
return nullptr;
}
if (auto* el = Type(t->type)) {
if (auto* vector = builder_->create<type::Vector>(el, t->width)) {
if (validator_.Vector(vector, t->source)) {
return vector;
}
}
}
return nullptr;
},
[&](const ast::Matrix* t) -> type::Matrix* {
if (!t->type) {
AddError("missing matrix element type", t->source.End());
return nullptr;
}
if (auto* el = Type(t->type)) {
if (auto* column_type = builder_->create<type::Vector>(el, t->rows)) {
if (auto* matrix = builder_->create<type::Matrix>(column_type, t->columns)) {
if (validator_.Matrix(matrix, t->source)) {
return matrix;
}
}
}
}
return nullptr;
},
[&](const ast::Array* t) { return Array(t); },
[&](const ast::Atomic* t) -> type::Atomic* {
if (auto* el = Type(t->type)) {
auto* a = builder_->create<type::Atomic>(el);
if (!validator_.Atomic(t, a)) {
return nullptr;
}
return a;
}
return nullptr;
},
[&](const ast::Pointer* t) -> type::Pointer* {
if (auto* el = Type(t->type)) {
auto access = t->access;
if (access == ast::Access::kUndefined) {
access = DefaultAccessForAddressSpace(t->address_space);
}
auto ptr = builder_->create<type::Pointer>(el, t->address_space, access);
if (!ptr) {
return nullptr;
}
if (!validator_.Pointer(t, ptr)) {
return nullptr;
}
if (!ApplyAddressSpaceUsageToType(t->address_space, el, t->type->source)) {
AddNote("while instantiating " + builder_->FriendlyName(ptr), t->source);
return nullptr;
}
return ptr;
}
return nullptr;
},
[&](const ast::Sampler* t) { return builder_->create<type::Sampler>(t->kind); },
[&](const ast::SampledTexture* t) -> type::SampledTexture* {
if (auto* el = Type(t->type)) {
auto* sem = builder_->create<type::SampledTexture>(t->dim, el);
if (!validator_.SampledTexture(sem, t->source)) {
return nullptr;
}
return sem;
}
return nullptr;
},
[&](const ast::MultisampledTexture* t) -> type::MultisampledTexture* {
if (auto* el = Type(t->type)) {
auto* sem = builder_->create<type::MultisampledTexture>(t->dim, el);
if (!validator_.MultisampledTexture(sem, t->source)) {
return nullptr;
}
return sem;
}
return nullptr;
},
[&](const ast::DepthTexture* t) { return builder_->create<type::DepthTexture>(t->dim); },
[&](const ast::DepthMultisampledTexture* t) {
return builder_->create<type::DepthMultisampledTexture>(t->dim);
},
[&](const ast::StorageTexture* t) -> type::StorageTexture* {
if (auto* el = Type(t->type)) {
if (!validator_.StorageTexture(t)) {
return nullptr;
}
return builder_->create<type::StorageTexture>(t->dim, t->format, t->access, el);
}
return nullptr;
},
[&](const ast::ExternalTexture*) { return builder_->create<type::ExternalTexture>(); },
[&](Default) {
auto* resolved = sem_.ResolvedSymbol(ty);
return Switch(
resolved, //
[&](type::Type* type) { return type; },
[&](sem::Variable* var) {
auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
AddError("cannot use variable '" + name + "' as type", ty->source);
AddNote("'" + name + "' declared here", var->Declaration()->source);
return nullptr;
},
[&](sem::Function* func) {
auto name = builder_->Symbols().NameFor(func->Declaration()->symbol);
AddError("cannot use function '" + name + "' as type", ty->source);
AddNote("'" + name + "' declared here", func->Declaration()->source);
return nullptr;
},
[&](Default) -> type::Type* {
if (auto* tn = ty->As<ast::TypeName>()) {
if (IsBuiltin(tn->name)) {
auto name = builder_->Symbols().NameFor(tn->name);
AddError("cannot use builtin '" + name + "' as type", ty->source);
return nullptr;
}
return ShortName(tn->name, tn->source);
}
TINT_UNREACHABLE(Resolver, diagnostics_)
<< "Unhandled resolved type '"
<< (resolved ? resolved->TypeInfo().name : "<null>")
<< "' resolved from ast::Type '" << ty->TypeInfo().name << "'";
return nullptr;
});
});
if (s) {
builder_->Sem().Add(ty, s);
}
return s;
}
sem::Variable* Resolver::Variable(const ast::Variable* v, bool is_global) {
return Switch(
v, //
[&](const ast::Var* var) { return Var(var, is_global); },
[&](const ast::Let* let) { return Let(let, is_global); },
[&](const ast::Override* override) { return Override(override); },
[&](const ast::Const* const_) { return Const(const_, is_global); },
[&](Default) {
TINT_ICE(Resolver, diagnostics_)
<< "Resolver::GlobalVariable() called with a unknown variable type: "
<< v->TypeInfo().name;
return nullptr;
});
}
sem::Variable* Resolver::Let(const ast::Let* v, bool is_global) {
const type::Type* ty = nullptr;
// If the variable has a declared type, resolve it.
if (v->type) {
ty = Type(v->type);
if (!ty) {
return nullptr;
}
}
if (!v->initializer) {
AddError("'let' declaration must have an initializer", v->source);
return nullptr;
}
auto* rhs = Load(Materialize(Expression(v->initializer), ty));
if (!rhs) {
return nullptr;
}
// If the variable has no declared type, infer it from the RHS
if (!ty) {
ty = rhs->Type()->UnwrapRef(); // Implicit load of RHS
}
if (rhs && !validator_.VariableInitializer(v, ast::AddressSpace::kNone, ty, rhs)) {
return nullptr;
}
if (!ApplyAddressSpaceUsageToType(ast::AddressSpace::kNone, const_cast<type::Type*>(ty),
v->source)) {
AddNote("while instantiating 'let' " + builder_->Symbols().NameFor(v->symbol), v->source);
return nullptr;
}
sem::Variable* sem = nullptr;
if (is_global) {
sem = builder_->create<sem::GlobalVariable>(
v, ty, sem::EvaluationStage::kRuntime, ast::AddressSpace::kNone,
ast::Access::kUndefined,
/* constant_value */ nullptr, sem::BindingPoint{}, std::nullopt);
} else {
sem = builder_->create<sem::LocalVariable>(v, ty, sem::EvaluationStage::kRuntime,
ast::AddressSpace::kNone,
ast::Access::kUndefined, current_statement_,
/* constant_value */ nullptr);
}
sem->SetInitializer(rhs);
builder_->Sem().Add(v, sem);
return sem;
}
sem::Variable* Resolver::Override(const ast::Override* v) {
const type::Type* ty = nullptr;
// If the variable has a declared type, resolve it.
if (v->type) {
ty = Type(v->type);
if (!ty) {
return nullptr;
}
}
const sem::Expression* rhs = nullptr;
// Does the variable have an initializer?
if (v->initializer) {
// Note: RHS must be a const or override expression, which excludes references.
// So there's no need to load or unwrap references here.
ExprEvalStageConstraint constraint{sem::EvaluationStage::kOverride, "override initializer"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
rhs = Materialize(Expression(v->initializer), ty);
if (!rhs) {
return nullptr;
}
// If the variable has no declared type, infer it from the RHS
if (!ty) {
ty = rhs->Type();
}
} else if (!ty) {
AddError("override declaration requires a type or initializer", v->source);
return nullptr;
}
if (rhs && !validator_.VariableInitializer(v, ast::AddressSpace::kNone, ty, rhs)) {
return nullptr;
}
if (!ApplyAddressSpaceUsageToType(ast::AddressSpace::kNone, const_cast<type::Type*>(ty),
v->source)) {
AddNote("while instantiating 'override' " + builder_->Symbols().NameFor(v->symbol),
v->source);
return nullptr;
}
auto* sem = builder_->create<sem::GlobalVariable>(
v, ty, sem::EvaluationStage::kOverride, ast::AddressSpace::kNone, ast::Access::kUndefined,
/* constant_value */ nullptr, sem::BindingPoint{}, std::nullopt);
sem->SetInitializer(rhs);
if (auto* id_attr = ast::GetAttribute<ast::IdAttribute>(v->attributes)) {
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@id"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* materialized = Materialize(Expression(id_attr->expr));
if (!materialized) {
return nullptr;
}
if (!materialized->Type()->IsAnyOf<type::I32, type::U32>()) {
AddError("@id must be an i32 or u32 value", id_attr->source);
return nullptr;
}
auto const_value = materialized->ConstantValue();
auto value = const_value->ValueAs<AInt>();
if (value < 0) {
AddError("@id value must be non-negative", id_attr->source);
return nullptr;
}
if (value > std::numeric_limits<decltype(OverrideId::value)>::max()) {
AddError("@id value must be between 0 and " +
std::to_string(std::numeric_limits<decltype(OverrideId::value)>::max()),
id_attr->source);
return nullptr;
}
auto o = OverrideId{static_cast<decltype(OverrideId::value)>(value)};
sem->SetOverrideId(o);
// Track the constant IDs that are specified in the shader.
override_ids_.Add(o, sem);
}
builder_->Sem().Add(v, sem);
return sem;
}
sem::Variable* Resolver::Const(const ast::Const* c, bool is_global) {
const type::Type* ty = nullptr;
// If the variable has a declared type, resolve it.
if (c->type) {
ty = Type(c->type);
if (!ty) {
return nullptr;
}
}
if (!c->initializer) {
AddError("'const' declaration must have an initializer", c->source);
return nullptr;
}
const sem::Expression* rhs = nullptr;
{
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "const initializer"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
rhs = Expression(c->initializer);
if (!rhs) {
return nullptr;
}
}
// Note: RHS must be a const expression, which excludes references.
// So there's no need to load or unwrap references here.
if (ty) {
// If an explicit type was specified, materialize to that type
rhs = Materialize(rhs, ty);
if (!rhs) {
return nullptr;
}
} else {
// If no type was specified, infer it from the RHS
ty = rhs->Type();
}
if (!validator_.VariableInitializer(c, ast::AddressSpace::kNone, ty, rhs)) {
return nullptr;
}
if (!ApplyAddressSpaceUsageToType(ast::AddressSpace::kNone, const_cast<type::Type*>(ty),
c->source)) {
AddNote("while instantiating 'const' " + builder_->Symbols().NameFor(c->symbol), c->source);
return nullptr;
}
const auto value = rhs->ConstantValue();
auto* sem = is_global ? static_cast<sem::Variable*>(builder_->create<sem::GlobalVariable>(
c, ty, sem::EvaluationStage::kConstant, ast::AddressSpace::kNone,
ast::Access::kUndefined, value, sem::BindingPoint{}, std::nullopt))
: static_cast<sem::Variable*>(builder_->create<sem::LocalVariable>(
c, ty, sem::EvaluationStage::kConstant, ast::AddressSpace::kNone,
ast::Access::kUndefined, current_statement_, value));
sem->SetInitializer(rhs);
builder_->Sem().Add(c, sem);
return sem;
}
sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
const type::Type* storage_ty = nullptr;
// If the variable has a declared type, resolve it.
if (auto* ty = var->type) {
storage_ty = Type(ty);
if (!storage_ty) {
return nullptr;
}
}
const sem::Expression* rhs = nullptr;
// Does the variable have a initializer?
if (var->initializer) {
ExprEvalStageConstraint constraint{
is_global ? sem::EvaluationStage::kOverride : sem::EvaluationStage::kRuntime,
"var initializer",
};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
rhs = Load(Materialize(Expression(var->initializer), storage_ty));
if (!rhs) {
return nullptr;
}
// If the variable has no declared type, infer it from the RHS
if (!storage_ty) {
storage_ty = rhs->Type();
}
}
if (!storage_ty) {
AddError("var declaration requires a type or initializer", var->source);
return nullptr;
}
auto address_space = var->declared_address_space;
if (address_space == ast::AddressSpace::kNone) {
// No declared address space. Infer from usage / type.
if (!is_global) {
address_space = ast::AddressSpace::kFunction;
} else if (storage_ty->UnwrapRef()->is_handle()) {
// https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
// If the store type is a texture type or a sampler type, then the
// variable declaration must not have a address space attribute. The
// address space will always be handle.
address_space = ast::AddressSpace::kHandle;
}
}
if (!is_global && address_space != ast::AddressSpace::kFunction &&
validator_.IsValidationEnabled(var->attributes,
ast::DisabledValidation::kIgnoreAddressSpace)) {
AddError("function-scope 'var' declaration must use 'function' address space", var->source);
return nullptr;
}
auto access = var->declared_access;
if (access == ast::Access::kUndefined) {
access = DefaultAccessForAddressSpace(address_space);
}
if (rhs && !validator_.VariableInitializer(var, address_space, storage_ty, rhs)) {
return nullptr;
}
auto* var_ty = builder_->create<type::Reference>(storage_ty, address_space, access);
if (!ApplyAddressSpaceUsageToType(address_space, var_ty,
var->type ? var->type->source : var->source)) {
AddNote("while instantiating 'var' " + builder_->Symbols().NameFor(var->symbol),
var->source);
return nullptr;
}
sem::Variable* sem = nullptr;
if (is_global) {
sem::BindingPoint binding_point;
if (var->HasBindingPoint()) {
uint32_t binding = 0;
{
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@binding"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* attr = ast::GetAttribute<ast::BindingAttribute>(var->attributes);
auto* materialized = Materialize(Expression(attr->expr));
if (!materialized) {
return nullptr;
}
if (!materialized->Type()->IsAnyOf<type::I32, type::U32>()) {
AddError("@binding must be an i32 or u32 value", attr->source);
return nullptr;
}
auto const_value = materialized->ConstantValue();
auto value = const_value->ValueAs<AInt>();
if (value < 0) {
AddError("@binding value must be non-negative", attr->source);
return nullptr;
}
binding = u32(value);
}
uint32_t group = 0;
{
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@group"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* attr = ast::GetAttribute<ast::GroupAttribute>(var->attributes);
auto* materialized = Materialize(Expression(attr->expr));
if (!materialized) {
return nullptr;
}
if (!materialized->Type()->IsAnyOf<type::I32, type::U32>()) {
AddError("@group must be an i32 or u32 value", attr->source);
return nullptr;
}
auto const_value = materialized->ConstantValue();
auto value = const_value->ValueAs<AInt>();
if (value < 0) {
AddError("@group value must be non-negative", attr->source);
return nullptr;
}
group = u32(value);
}
binding_point = {group, binding};
}
std::optional<uint32_t> location;
if (auto* attr = ast::GetAttribute<ast::LocationAttribute>(var->attributes)) {
auto value = LocationAttribute(attr);
if (!value) {
return nullptr;
}
location = value.Get();
}
sem = builder_->create<sem::GlobalVariable>(
var, var_ty, sem::EvaluationStage::kRuntime, address_space, access,
/* constant_value */ nullptr, binding_point, location);
} else {
sem = builder_->create<sem::LocalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
address_space, access, current_statement_,
/* constant_value */ nullptr);
}
sem->SetInitializer(rhs);
builder_->Sem().Add(var, sem);
return sem;
}
sem::Parameter* Resolver::Parameter(const ast::Parameter* param, uint32_t index) {
auto add_note = [&] {
AddNote("while instantiating parameter " + builder_->Symbols().NameFor(param->symbol),
param->source);
};
for (auto* attr : param->attributes) {
Mark(attr);
}
if (!validator_.NoDuplicateAttributes(param->attributes)) {
return nullptr;
}
type::Type* ty = Type(param->type);
if (!ty) {
return nullptr;
}
if (!ApplyAddressSpaceUsageToType(ast::AddressSpace::kNone, ty, param->type->source)) {
add_note();
return nullptr;
}
if (auto* ptr = ty->As<type::Pointer>()) {
// For MSL, we push module-scope variables into the entry point as pointer
// parameters, so we also need to handle their store type.
if (!ApplyAddressSpaceUsageToType(
ptr->AddressSpace(), const_cast<type::Type*>(ptr->StoreType()), param->source)) {
add_note();
return nullptr;
}
}
sem::BindingPoint binding_point;
if (param->HasBindingPoint()) {
{
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@binding value"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* attr = ast::GetAttribute<ast::BindingAttribute>(param->attributes);
auto* materialized = Materialize(Expression(attr->expr));
if (!materialized) {
return nullptr;
}
binding_point.binding = materialized->ConstantValue()->ValueAs<u32>();
}
{
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@group value"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* attr = ast::GetAttribute<ast::GroupAttribute>(param->attributes);
auto* materialized = Materialize(Expression(attr->expr));
if (!materialized) {
return nullptr;
}
binding_point.group = materialized->ConstantValue()->ValueAs<u32>();
}
}
std::optional<uint32_t> location;
if (auto* attr = ast::GetAttribute<ast::LocationAttribute>(param->attributes)) {
auto value = LocationAttribute(attr);
if (!value) {
return nullptr;
}
location = value.Get();
}
auto* sem = builder_->create<sem::Parameter>(
param, index, ty, ast::AddressSpace::kNone, ast::Access::kUndefined,
sem::ParameterUsage::kNone, binding_point, location);
builder_->Sem().Add(param, sem);
return sem;
}
utils::Result<uint32_t> Resolver::LocationAttribute(const ast::LocationAttribute* attr) {
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "@location value"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* materialized = Materialize(Expression(attr->expr));
if (!materialized) {
return utils::Failure;
}
if (!materialized->Type()->IsAnyOf<type::I32, type::U32>()) {
AddError("@location must be an i32 or u32 value", attr->source);
return utils::Failure;
}
auto const_value = materialized->ConstantValue();
auto value = const_value->ValueAs<AInt>();
if (value < 0) {
AddError("@location value must be non-negative", attr->source);
return utils::Failure;
}
return static_cast<uint32_t>(value);
}
ast::Access Resolver::DefaultAccessForAddressSpace(ast::AddressSpace address_space) {
// https://gpuweb.github.io/gpuweb/wgsl/#storage-class
switch (address_space) {
case ast::AddressSpace::kStorage:
case ast::AddressSpace::kUniform:
case ast::AddressSpace::kHandle:
return ast::Access::kRead;
default:
break;
}
return ast::Access::kReadWrite;
}
bool Resolver::AllocateOverridableConstantIds() {
constexpr size_t kLimit = std::numeric_limits<decltype(OverrideId::value)>::max();
// The next pipeline constant ID to try to allocate.
OverrideId next_id;
bool ids_exhausted = false;
auto increment_next_id = [&] {
if (next_id.value == kLimit) {
ids_exhausted = true;
} else {
next_id.value = next_id.value + 1;
}
};
// Allocate constant IDs in global declaration order, so that they are
// deterministic.
// TODO(crbug.com/tint/1192): If a transform changes the order or removes an
// unused constant, the allocation may change on the next Resolver pass.
for (auto* decl : builder_->AST().GlobalDeclarations()) {
auto* override = decl->As<ast::Override>();
if (!override) {
continue;
}
OverrideId id;
if (ast::HasAttribute<ast::IdAttribute>(override->attributes)) {
id = builder_->Sem().Get<sem::GlobalVariable>(override)->OverrideId();
} else {
// No ID was specified, so allocate the next available ID.
while (!ids_exhausted && override_ids_.Contains(next_id)) {
increment_next_id();
}
if (ids_exhausted) {
AddError(
"number of 'override' variables exceeded limit of " + std::to_string(kLimit),
decl->source);
return false;
}
id = next_id;
increment_next_id();
}
auto* sem = sem_.Get<sem::GlobalVariable>(override);
const_cast<sem::GlobalVariable*>(sem)->SetOverrideId(id);
}
return true;
}
void Resolver::SetShadows() {
for (auto it : dependencies_.shadows) {
CastableBase* b = sem_.Get(it.value);
if (!b) {
TINT_ICE(Resolver, builder_->Diagnostics())
<< "AST node '" << it.value->TypeInfo().name << "' had no semantic info\n"
<< "At: " << it.value->source << "\n"
<< "Pointer: " << it.value;
}
Switch(
sem_.Get(it.key), //
[&](sem::LocalVariable* local) { local->SetShadows(b); },
[&](sem::Parameter* param) { param->SetShadows(b); });
}
}
sem::GlobalVariable* Resolver::GlobalVariable(const ast::Variable* v) {
utils::UniqueVector<const sem::GlobalVariable*, 4> transitively_referenced_overrides;
TINT_SCOPED_ASSIGNMENT(resolved_overrides_, &transitively_referenced_overrides);
auto* sem = As<sem::GlobalVariable>(Variable(v, /* is_global */ true));
if (!sem) {
return nullptr;
}
for (auto* attr : v->attributes) {
Mark(attr);
}
if (!validator_.NoDuplicateAttributes(v->attributes)) {
return nullptr;
}
if (!validator_.GlobalVariable(sem, override_ids_)) {
return nullptr;
}
// Track the pipeline-overridable constants that are transitively referenced by this variable.
for (auto* var : transitively_referenced_overrides) {
builder_->Sem().AddTransitivelyReferencedOverride(sem, var);
}
if (auto* arr = sem->Type()->UnwrapRef()->As<type::Array>()) {
auto* refs = builder_->Sem().TransitivelyReferencedOverrides(arr);
if (refs) {
for (auto* var : *refs) {
builder_->Sem().AddTransitivelyReferencedOverride(sem, var);
}
}
}
return sem;
}
sem::Statement* Resolver::StaticAssert(const ast::StaticAssert* assertion) {
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "static assertion"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
auto* expr = Expression(assertion->condition);
if (!expr) {
return nullptr;
}
auto* cond = expr->ConstantValue();
if (auto* ty = cond->Type(); !ty->Is<type::Bool>()) {
AddError(
"static assertion condition must be a bool, got '" + builder_->FriendlyName(ty) + "'",
assertion->condition->source);
return nullptr;
}
if (!cond->ValueAs<bool>()) {
AddError("static assertion failed", assertion->source);
return nullptr;
}
auto* sem =
builder_->create<sem::Statement>(assertion, current_compound_statement_, current_function_);
builder_->Sem().Add(assertion, sem);
return sem;
}
sem::Function* Resolver::Function(const ast::Function* decl) {
uint32_t parameter_index = 0;
utils::Hashmap<Symbol, Source, 8> parameter_names;
utils::Vector<sem::Parameter*, 8> parameters;
// Resolve all the parameters
for (auto* param : decl->params) {
Mark(param);
{ // Check the parameter name is unique for the function
if (auto added = parameter_names.Add(param->symbol, param->source); !added) {
auto name = builder_->Symbols().NameFor(param->symbol);
AddError("redefinition of parameter '" + name + "'", param->source);
AddNote("previous definition is here", *added.value);
return nullptr;
}
}
auto* p = Parameter(param, parameter_index++);
if (!p) {
return nullptr;
}
if (!validator_.Parameter(decl, p)) {
return nullptr;
}
parameters.Push(p);
auto* p_ty = const_cast<type::Type*>(p->Type());
if (auto* str = p_ty->As<sem::Struct>()) {
switch (decl->PipelineStage()) {
case ast::PipelineStage::kVertex:
str->AddUsage(type::PipelineStageUsage::kVertexInput);
break;
case ast::PipelineStage::kFragment:
str->AddUsage(type::PipelineStageUsage::kFragmentInput);
break;
case ast::PipelineStage::kCompute:
str->AddUsage(type::PipelineStageUsage::kComputeInput);
break;
case ast::PipelineStage::kNone:
break;
}
}
}
// Resolve the return type
type::Type* return_type = nullptr;
if (auto* ty = decl->return_type) {
return_type = Type(ty);
if (!return_type) {
return nullptr;
}
} else {
return_type = builder_->create<type::Void>();
}
// Determine if the return type has a location
std::optional<uint32_t> return_location;
for (auto* attr : decl->return_type_attributes) {
Mark(attr);
if (auto* loc_attr = attr->As<ast::LocationAttribute>()) {
auto value = LocationAttribute(loc_attr);
if (!value) {
return nullptr;
}
return_location = value.Get();
}
}
if (!validator_.NoDuplicateAttributes(decl->attributes)) {
return nullptr;
}
if (auto* str = return_type->As<sem::Struct>()) {
if (!ApplyAddressSpaceUsageToType(ast::AddressSpace::kNone, str, decl->source)) {
AddNote(
"while instantiating return type for " + builder_->Symbols().NameFor(decl->symbol),
decl->source);
return nullptr;
}
switch (decl->PipelineStage()) {
case ast::PipelineStage::kVertex:
str->AddUsage(type::PipelineStageUsage::kVertexOutput);
break;
case ast::PipelineStage::kFragment:
str->AddUsage(type::PipelineStageUsage::kFragmentOutput);
break;
case ast::PipelineStage::kCompute:
str->AddUsage(type::PipelineStageUsage::kComputeOutput);
break;
case ast::PipelineStage::kNone:
break;
}
}
auto* func =
builder_->create<sem::Function>(decl, return_type, return_location, std::move(parameters));
builder_->Sem().Add(decl, func);
TINT_SCOPED_ASSIGNMENT(current_function_, func);
if (!WorkgroupSize(decl)) {
return nullptr;
}
if (decl->IsEntryPoint()) {
entry_points_.Push(func);
}
if (decl->body) {
Mark(decl->body);
if (current_compound_statement_) {
TINT_ICE(Resolver, diagnostics_)
<< "Resolver::Function() called with a current compound statement";
return nullptr;
}
auto* body = StatementScope(decl->body, builder_->create<sem::FunctionBlockStatement>(func),
[&] { return Statements(decl->body->statements); });
if (!body) {
return nullptr;
}
func->Behaviors() = body->Behaviors();
if (func->Behaviors().Contains(sem::Behavior::kReturn)) {
// https://www.w3.org/TR/WGSL/#behaviors-rules
// We assign a behavior to each function: it is its body’s behavior
// (treating the body as a regular statement), with any "Return" replaced
// by "Next".
func->Behaviors().Remove(sem::Behavior::kReturn);
func->Behaviors().Add(sem::Behavior::kNext);
}
}
for (auto* attr : decl->attributes) {
Mark(attr);
}
if (!validator_.NoDuplicateAttributes(decl->return_type_attributes)) {
return nullptr;
}
auto stage = current_function_ ? current_function_->Declaration()->PipelineStage()
: ast::PipelineStage::kNone;
if (!validator_.Function(func, stage)) {
return nullptr;
}
// If this is an entry point, mark all transitively called functions as being
// used by this entry point.
if (decl->IsEntryPoint()) {
for (auto* f : func->TransitivelyCalledFunctions()) {
const_cast<sem::Function*>(f)->AddAncestorEntryPoint(func);
}
}
return func;
}
bool Resolver::WorkgroupSize(const ast::Function* func) {
// Set work-group size defaults.
sem::WorkgroupSize ws;
for (size_t i = 0; i < 3; i++) {
ws[i] = 1;
}
auto* attr = ast::GetAttribute<ast::WorkgroupAttribute>(func->attributes);
if (!attr) {
return true;
}
auto values = attr->Values();
utils::Vector<const sem::Expression*, 3> args;
utils::Vector<const type::Type*, 3> arg_tys;
constexpr const char* kErrBadExpr =
"workgroup_size argument must be a constant or override-expression of type "
"abstract-integer, i32 or u32";
for (size_t i = 0; i < 3; i++) {
// Each argument to this attribute can either be a literal, an identifier for a module-scope
// constants, a const-expression, or nullptr if not specified.
auto* value = values[i];
if (!value) {
break;
}
const auto* expr = Expression(value);
if (!expr) {
return false;
}
auto* ty = expr->Type();
if (!ty->IsAnyOf<type::I32, type::U32, type::AbstractInt>()) {
AddError(kErrBadExpr, value->source);
return false;
}
if (expr->Stage() != sem::EvaluationStage::kConstant &&
expr->Stage() != sem::EvaluationStage::kOverride) {
AddError(kErrBadExpr, value->source);
return false;
}
args.Push(expr);
arg_tys.Push(ty);
}
auto* common_ty = type::Type::Common(arg_tys);
if (!common_ty) {
AddError("workgroup_size arguments must be of the same type, either i32 or u32",
attr->source);
return false;
}
// If all arguments are abstract-integers, then materialize to i32.
if (common_ty->Is<type::AbstractInt>()) {
common_ty = builder_->create<type::I32>();
}
for (size_t i = 0; i < args.Length(); i++) {
auto* materialized = Materialize(args[i], common_ty);
if (!materialized) {
return false;
}
if (auto* value = materialized->ConstantValue()) {
if (value->ValueAs<AInt>() < 1) {
AddError("workgroup_size argument must be at least 1", values[i]->source);
return false;
}
ws[i] = value->ValueAs<u32>();
} else {
ws[i] = std::nullopt;
}
}
uint64_t total_size = static_cast<uint64_t>(ws[0].value_or(1));
for (size_t i = 1; i < 3; i++) {
total_size *= static_cast<uint64_t>(ws[i].value_or(1));
if (total_size > 0xffffffff) {
AddError("total workgroup grid size cannot exceed 0xffffffff", values[i]->source);
return false;
}
}
current_function_->SetWorkgroupSize(std::move(ws));
return true;
}
bool Resolver::Statements(utils::VectorRef<const ast::Statement*> stmts) {
sem::Behaviors behaviors{sem::Behavior::kNext};
bool reachable = true;
for (auto* stmt : stmts) {
Mark(stmt);
auto* sem = Statement(stmt);
if (!sem) {
return false;
}
// s1 s2:(B1∖{Next}) ∪ B2
sem->SetIsReachable(reachable);
if (reachable) {
behaviors = (behaviors - sem::Behavior::kNext) + sem->Behaviors();
}
reachable = reachable && sem->Behaviors().Contains(sem::Behavior::kNext);
}
current_statement_->Behaviors() = behaviors;
if (!validator_.Statements(stmts)) {
return false;
}
return true;
}
sem::Statement* Resolver::Statement(const ast::Statement* stmt) {
return Switch(
stmt,
// Compound statements. These create their own sem::CompoundStatement
// bindings.
[&](const ast::BlockStatement* b) { return BlockStatement(b); },
[&](const ast::ForLoopStatement* l) { return ForLoopStatement(l); },
[&](const ast::LoopStatement* l) { return LoopStatement(l); },
[&](const ast::WhileStatement* w) { return WhileStatement(w); },
[&](const ast::IfStatement* i) { return IfStatement(i); },
[&](const ast::SwitchStatement* s) { return SwitchStatement(s); },
// Non-Compound statements
[&](const ast::AssignmentStatement* a) { return AssignmentStatement(a); },
[&](const ast::BreakStatement* b) { return BreakStatement(b); },
[&](const ast::BreakIfStatement* b) { return BreakIfStatement(b); },
[&](const ast::CallStatement* c) { return CallStatement(c); },
[&](const ast::CompoundAssignmentStatement* c) { return CompoundAssignmentStatement(c); },
[&](const ast::ContinueStatement* c) { return ContinueStatement(c); },
[&](const ast::DiscardStatement* d) { return DiscardStatement(d); },
[&](const ast::IncrementDecrementStatement* i) { return IncrementDecrementStatement(i); },
[&](const ast::ReturnStatement* r) { return ReturnStatement(r); },
[&](const ast::VariableDeclStatement* v) { return VariableDeclStatement(v); },
[&](const ast::StaticAssert* sa) { return StaticAssert(sa); },
// Error cases
[&](const ast::CaseStatement*) {
AddError("case statement can only be used inside a switch statement", stmt->source);
return nullptr;
},
[&](Default) {
AddError("unknown statement type: " + std::string(stmt->TypeInfo().name), stmt->source);
return nullptr;
});
}
sem::CaseStatement* Resolver::CaseStatement(const ast::CaseStatement* stmt, const type::Type* ty) {
auto* sem =
builder_->create<sem::CaseStatement>(stmt, current_compound_statement_, current_function_);
return StatementScope(stmt, sem, [&] {
sem->Selectors().reserve(stmt->selectors.Length());
for (auto* sel : stmt->selectors) {
Mark(sel);
ExprEvalStageConstraint constraint{sem::EvaluationStage::kConstant, "case selector"};
TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
const constant::Value* const_value = nullptr;
if (!sel->IsDefault()) {
// The sem statement was created in the switch when attempting to determine the
// common type.
auto* materialized = Materialize(sem_.Get(sel->expr), ty);
if (!materialized) {
return false;
}
if (!materialized->Type()->IsAnyOf<type::I32, type::U32>()) {
AddError("case selector must be an i32 or u32 value", sel->source);
return false;
}
const_value = materialized->ConstantValue();
if (!const_value) {
AddError("case selector must be a constant expression", sel->source);
return false;
}
}
sem->Selectors().emplace_back(builder_->create<sem::CaseSelector>(sel, const_value));
}
Mark(stmt->body);
auto* body = BlockStatement(stmt->body);
if (!body) {
return false;
}
sem->SetBlock(body);
sem->Behaviors() = body->Behaviors();
return true;
});
}
sem::IfStatement* Resolver::IfStatement(const ast::IfStatement* stmt) {
auto* sem =
builder_->create<sem::IfStatement>(stmt, current_compound_statement_, current_function_);
return StatementScope(stmt, sem, [&] {
auto* cond = Load(Expression(stmt->condition));
if (!cond) {
return false;
}
sem->SetCondition(cond);
sem->Behaviors() = cond->Behaviors();
sem->Behaviors().Remove(sem::Behavior::kNext);
Mark(stmt->body);
auto* body = builder_->create<sem::BlockStatement>(stmt->body, current_compound_statement_,
current_function_);
if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
return false;
}
sem->Behaviors().Add(body->Behaviors());
if (stmt->else_statement) {
Mark(stmt->else_statement);
auto* else_sem = Statement(stmt->else_statement);
if (!else_sem) {
return false;
}
sem->Behaviors().Add(else_sem->Behaviors());
} else {
// https://www.w3.org/TR/WGSL/#behaviors-rules
// if statements without an else branch are treated as if they had an
// empty else branch (which adds Next to their behavior)
sem->Behaviors().Add(sem::Behavior::kNext);
}
return validator_.IfStatement(sem);
});
}
sem::BlockStatement* Resolver::BlockStatement(const ast::BlockStatement* stmt) {
auto* sem = builder_->create<sem::BlockStatement>(
stmt->As<ast::BlockStatement>(), current_compound_statement_, current_function_);
return StatementScope(stmt, sem, [&] { return Statements(stmt->statements); });
}
sem::LoopStatement* Resolver::LoopStatement(const ast::LoopStatement* stmt) {
auto* sem =
builder_->create<sem::LoopStatement>(stmt, current_compound_statement_, current_function_);
return StatementScope(stmt, sem, [&] {
Mark(stmt->body);
auto* body = builder_->create<sem::LoopBlockStatement>(
stmt->body, current_compound_statement_, current_function_);
return StatementScope(stmt->body, body, [&] {
if (!Statements(stmt->body->statements)) {
return false;
}
auto& behaviors = sem->Behaviors();
behaviors = body->Behaviors();
if (stmt->continuing) {
Mark(stmt->continuing);
auto* continuing = StatementScope(
stmt->continuing,
builder_->create<sem::LoopContinuingBlockStatement>(
stmt->continuing, current_compound_statement_, current_function_),
[&] { return Statements(stmt->continuing->statements); });
if (!continuing) {
return false;
}
behaviors.Add(continuing->Behaviors());
}
if (behaviors.Contains(sem::Behavior::kBreak)) { // Does the loop exit?
behaviors.Add(sem::Behavior::kNext);
} else {
behaviors.Remove(sem::Behavior::kNext);
}
behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
return validator_.LoopStatement(sem);
});
});
}
sem::ForLoopStatement* Resolver::ForLoopStatement(const ast::ForLoopStatement* stmt) {
auto* sem = builder_->create<sem::ForLoopStatement>(stmt, current_compound_statement_,
current_function_);
return StatementScope(stmt, sem, [&] {
auto& behaviors = sem->Behaviors();
if (auto* initializer = stmt->initializer) {
Mark(initializer);
auto* init = Statement(initializer);
if (!init) {
return false;
}
behaviors.Add(init->Behaviors());
}
if (auto* cond_expr = stmt->condition) {
auto* cond = Load(Expression(cond_expr));
if (!cond) {
return false;
}
sem->SetCondition(cond);
behaviors.Add(cond->Behaviors());
}
if (auto* continuing = stmt->continuing) {
Mark(continuing);
auto* cont = Statement(continuing);
if (!cont) {
return false;
}
behaviors.Add(cont->Behaviors());
}
Mark(stmt->body);
auto* body = builder_->create<sem::LoopBlockStatement>(
stmt->body, current_compound_statement_, current_function_);
if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
return false;
}
behaviors.Add(body->Behaviors());
if (stmt->condition || behaviors.Contains(sem::Behavior::kBreak)) { // Does the loop exit?
behaviors.Add(sem::Behavior::kNext);
} else {
behaviors.Remove(sem::Behavior::kNext);
}
behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
return validator_.ForLoopStatement(sem);
});
}
sem::WhileStatement* Resolver::WhileStatement(const ast::WhileStatement* stmt) {
auto* sem =
builder_->create<sem::WhileStatement>(stmt, current_compound_statement_, current_function_);
return StatementScope(stmt, sem, [&] {
auto& behaviors = sem->Behaviors();
auto* cond = Load(Expression(stmt->condition));
if (!cond) {
return false;
}
sem->SetCondition(cond);
behaviors.Add(cond->Behaviors());
Mark(stmt->body);
auto* body = builder_->create<sem::LoopBlockStatement>(
stmt->body, current_compound_statement_, current_function_);
if (!StatementScope(stmt->body, body, [&] { return Statements(stmt->body->statements); })) {
return false;
}
behaviors.Add(body->Behaviors());
// Always consider the while as having a 'next' behaviour because it has
// a condition. We don't check if the condition will terminate but it isn't
// valid to have an infinite loop in a WGSL program, so a non-terminating
// condition is already an invalid program.
behaviors.Add(sem::Behavior::kNext);
behaviors.Remove(sem::Behavior::kBreak, sem::Behavior::kContinue);
return validator_.WhileStatement(sem);
});
}
sem::Expression* Resolver::Expression(const ast::Expression* root) {
utils::Vector<const ast::Expression*, 64> sorted;
constexpr size_t kMaxExpressionDepth = 512U;
bool failed = false;
if (!ast::TraverseExpressions<ast::TraverseOrder::RightToLeft>(
root, diagnostics_, [&](const ast::Expression* expr, size_t depth) {
if (depth > kMaxExpressionDepth) {
AddError(
"reached max expression depth of " + std::to_string(kMaxExpressionDepth),
expr->source);
failed = true;
return ast::TraverseAction::Stop;
}
if (!Mark(expr)) {
failed = true;
return ast::TraverseAction::Stop;
}
if (auto* binary = expr->As<ast::BinaryExpression>();
binary && binary->IsLogical()) {
// Store potential const-eval short-circuit pair
logical_binary_lhs_to_parent_.Add(binary->lhs, binary);
}
sorted.Push(expr);
return ast::TraverseAction::Descend;
})) {
return nullptr;
}
if (failed) {
return nullptr;
}
for (auto* expr : utils::Reverse(sorted)) {
auto* sem_expr = Switch(
expr,
[&](const ast::IndexAccessorExpression* array) -> sem::Expression* {
return IndexAccessor(array);
},
[&](const ast::BinaryExpression* bin_op) -> sem::Expression* { return Binary(bin_op); },
[&](const ast::BitcastExpression* bitcast) -> sem::Expression* {
return Bitcast(bitcast);
},
[&](const ast::CallExpression* call) -> sem::Expression* { return Call(call); },
[&](const ast::IdentifierExpression* ident) -> sem::Expression* {
return Identifier(ident);
},
[&](const ast::LiteralExpression* literal) -> sem::Expression* {
return Literal(literal);
},
[&](const ast::MemberAccessorExpression* member) -> sem::Expression* {
return MemberAccessor(member);
},
[&](const ast::UnaryOpExpression* unary) -> sem::Expression* { return UnaryOp(unary); },
[&](const ast::PhonyExpression*) -> sem::Expression* {
return builder_->create<sem::Expression>(expr, builder_->create<type::Void>(),
sem::EvaluationStage::kRuntime,
current_statement_,
/* constant_value */ nullptr,
/* has_side_effects */ false);
},
[&](Default) {
TINT_ICE(Resolver, diagnostics_)
<< "unhandled expression type: " << expr->TypeInfo().name;
return nullptr;
});
if (!sem_expr) {
return nullptr;
}
if (auto* constraint = expr_eval_stage_constraint_.constraint) {
if (!validator_.EvaluationStage(sem_expr, expr_eval_stage_constraint_.stage,
constraint)) {
return nullptr;
}
}
builder_->Sem().Add(expr, sem_expr);
if (expr == root) {
return sem_expr;
}
// If we just processed the lhs of a constexpr logical binary expression, mark the rhs for
// short-circuiting.
if (sem_expr->ConstantValue()) {
if (auto binary = logical_binary_lhs_to_parent_.Find(expr)) {
const bool lhs_is_true = sem_expr->ConstantValue()->ValueAs<bool>();
if (((*binary)->IsLogicalAnd() && !lhs_is_true) ||
((*binary)->IsLogicalOr() && lhs_is_true)) {
// Mark entire expression tree to not const-evaluate
auto r = ast::TraverseExpressions( //
(*binary)->rhs, diagnostics_, [&](const ast::Expression* e) {
skip_const_eval_.Add(e);
return ast::TraverseAction::Descend;
});
if (!r) {
return nullptr;
}
}
}
}
}
TINT_ICE(Resolver, diagnostics_) << "Expression() did not find root node";
return nullptr;
}
void Resolver::RegisterStore(const sem::Expression* expr) {
auto& info = alias_analysis_infos_[current_function_];
Switch(
expr->RootIdentifier(),
[&](const sem::GlobalVariable* global) {
info.module_scope_writes.insert({global, expr});
},
[&](const sem::Parameter* param) { info.parameter_writes.insert(param); });
}
bool Resolver::AliasAnalysis(const sem::Call* call) {
auto* target = call->Target()->As<sem::Function>();
if (!target) {
return true;
}
if (validator_.IsValidationDisabled(target->Declaration()->attributes,
ast::DisabledValidation::kIgnorePointerAliasing)) {
return true;
}
// Helper to generate an aliasing error diagnostic.
struct Alias {
const sem::Expression* expr; // the "other expression"
enum { Argument, ModuleScope } type; // the type of the "other" expression
std::string access; // the access performed for the "other" expression
};
auto make_error = [&](const sem::Expression* arg, Alias&& var) {
// TODO(crbug.com/tint/1675): Switch to error and return false after deprecation period.
AddWarning("invalid aliased pointer argument", arg->Declaration()->source);
switch (var.type) {
case Alias::Argument:
AddNote("aliases with another argument passed here",
var.expr->Declaration()->source);
break;
case Alias::ModuleScope: {
auto* func = var.expr->Stmt()->Function();
auto func_name = builder_->Symbols().NameFor(func->Declaration()->symbol);
AddNote(
"aliases with module-scope variable " + var.access + " in '" + func_name + "'",
var.expr->Declaration()->source);
break;
}
}
return true;
};
auto& args = call->Arguments();
auto& target_info = alias_analysis_infos_[target];
auto& caller_info = alias_analysis_infos_[current_function_];
// Track the set of root identifiers that are read and written by arguments passed in this call.
std::unordered_map<const sem::Variable*, const sem::Expression*> arg_reads;
std::unordered_map<const sem::Variable*, const sem::Expression*> arg_writes;
for (size_t i = 0; i < args.Length(); i++) {
auto* arg = args[i];
if (!arg->Type()->Is<type::Pointer>()) {
continue;
}
auto* root = arg->RootIdentifier();
if (target_info.parameter_writes.count(target->Parameters()[i])) {
// Arguments that are written to can alias with any other argument or module-scope
// variable access.
if (arg_writes.count(root)) {
return make_error(arg, {arg_writes.at(root), Alias::Argument, "write"});
}
if (arg_reads.count(root)) {
return make_error(arg, {arg_reads.at(root), Alias::Argument, "read"});
}
if (target_info.module_scope_reads.count(root)) {
return make_error(
arg, {target_info.module_scope_reads.at(root), Alias::ModuleScope, "read"});
}
if (target_info.module_scope_writes.count(root)) {
return make_error(
arg, {target_info.module_scope_writes.at(root), Alias::ModuleScope, "write"});
}
arg_writes.insert({root, arg});
// Propagate the write access to the caller.
Switch(
root,
[&](const sem::GlobalVariable* global) {
caller_info.module_scope_writes.insert({global, arg});
},
[&](const sem::Parameter* param) { caller_info.parameter_writes.insert(param); });
} else if (target_info.parameter_reads.count(target->Parameters()[i])) {
// Arguments that are read from can alias with arguments or module-scope variables that
// are written to.
if (arg_writes.count(root)) {
return make_error(arg, {arg_writes.at(root), Alias::Argument, "write"});
}
if (target_info.module_scope_writes.count(root)) {
return make_error(
arg, {target_info.module_scope_writes.at(root), Alias::ModuleScope, "write"});
}
arg_reads.insert({root, arg});
// Propagate the read access to the caller.
Switch(
root,
[&](const sem::GlobalVariable* global) {
caller_info.module_scope_reads.insert({global, arg});
},
[&](const sem::Parameter* param) { caller_info.parameter_reads.insert(param); });
}
}
// Propagate module-scope variable uses to the caller.
for (auto read : target_info.module_scope_reads) {
caller_info.module_scope_reads.insert({read.first, read.second});
}
for (auto write : target_info.module_scope_writes) {
caller_info.module_scope_writes.insert({write.first, write.second});
}
return true;
}
const type::Type* Resolver::ConcreteType(const type::Type* ty,
const type::Type* target_ty,
const Source& source) {
auto i32 = [&] { return builder_->create<type::I32>(); };
auto f32 = [&] { return builder_->create<type::F32>(); };
auto i32v = [&](uint32_t width) { return builder_->create<type::Vector>(i32(), width); };
auto f32v = [&](uint32_t width) { return builder_->create<type::Vector>(f32(), width); };
auto f32m = [&](uint32_t columns, uint32_t rows) {
return builder_->create<type::Matrix>(f32v(rows), columns);
};
return Switch(
ty, //
[&](const type::AbstractInt*) { return target_ty ? target_ty : i32(); },
[&](const type::AbstractFloat*) { return target_ty ? target_ty : f32(); },
[&](const type::Vector* v) {
return Switch(
v->type(), //
[&](const type::AbstractInt*) { return target_ty ? target_ty : i32v(v->Width()); },
[&](const type::AbstractFloat*) {
return target_ty ? target_ty : f32v(v->Width());
});
},
[&](const type::Matrix* m) {
return Switch(m->type(), //
[&](const type::AbstractFloat*) {
return target_ty ? target_ty : f32m(m->columns(), m->rows());
});
},
[&](const type::Array* a) -> const type::Type* {
const type::Type* target_el_ty = nullptr;
if (auto* target_arr_ty = As<type::Array>(target_ty)) {
target_el_ty = target_arr_ty->ElemType();
}
if (auto* el_ty = ConcreteType(a->ElemType(), target_el_ty, source)) {
return Array(source, source, el_ty, a->Count(), /* explicit_stride */ 0);
}
return nullptr;
},
[&](const sem::Struct* s) -> const type::Type* {
if (auto tys = s->ConcreteTypes(); !tys.IsEmpty()) {
return target_ty ? target_ty : tys[0];
}
return nullptr;
});
}
const sem::Expression* Resolver::Load(const sem::Expression* expr) {
if (!expr) {
// Allow for Load(Expression(blah)), where failures pass through Load()
return nullptr;
}
if (!expr->Type()->Is<type::Reference>()) {
// Expression is not a reference type, so cannot be loaded. Just return expr.
return expr;
}
auto* load = builder_->create<sem::Load>(expr, current_statement_);
load->Behaviors() = expr->Behaviors();
builder_->Sem().Replace(expr->Declaration(), load);
// Track the load for the alias analysis.
auto& alias_info = alias_analysis_infos_[current_function_];
Switch(
expr->RootIdentifier(),
[&](const sem::GlobalVariable* global) {
alias_info.module_scope_reads.insert({global, expr});
},
[&](const sem::Parameter* param) { alias_info.parameter_reads.insert(param); });
return load;
}
const sem::Expression* Resolver::Materialize(const sem::Expression* expr,
const type::Type* target_type /* = nullptr */) {
if (!expr) {
// Allow for Materialize(Expression(blah)), where failures pass through Materialize()
return nullptr;
}
auto* decl = expr->Declaration();
auto* concrete_ty = ConcreteType(expr->Type(), target_type, decl->source);
if (!concrete_ty) {
return expr; // Does not require materialization
}
auto* src_ty = expr->Type();
if (!validator_.Materialize(concrete_ty, src_ty, decl->source)) {
return nullptr;
}
const constant::Value* materialized_val = nullptr;
if (!skip_const_eval_.Contains(decl)) {
auto expr_val = expr->ConstantValue();
if (!expr_val) {
TINT_ICE(Resolver, builder_->Diagnostics())
<< decl->source << "Materialize(" << decl->TypeInfo().name
<< ") called on expression with no constant value";
return nullptr;
}
auto val = const_eval_.Convert(concrete_ty, expr_val, decl->source);
if (!val) {
// Convert() has already failed and raised an diagnostic error.
return nullptr;
}
materialized_val = val.Get();
if (!materialized_val) {
TINT_ICE(Resolver, builder_->Diagnostics())
<< decl->source << "ConvertValue(" << builder_->FriendlyName(expr_val->Type())
<< " -> " << builder_->FriendlyName(concrete_ty) << ") returned invalid value";
return nullptr;
}
}
auto* m =
builder_->create<sem::Materialize>(expr, current_statement_, concrete_ty, materialized_val);
m->Behaviors() = expr->Behaviors();
builder_->Sem().Replace(decl, m);
return m;
}
template <size_t N>
bool Resolver::MaybeMaterializeAndLoadArguments(utils::Vector<const sem::Expression*, N>& args,
const sem::CallTarget* target) {
for (size_t i = 0, n = std::min(args.Length(), target->Parameters().Length()); i < n; i++) {
const auto* param_ty = target->Parameters()[i]->Type();
if (ShouldMaterializeArgument(param_ty)) {
auto* materialized = Materialize(args[i], param_ty);
if (!materialized) {
return false;
}
args[i] = materialized;
}
if (!param_ty->Is<type::Reference>()) {
auto* load = Load(args[i]);
if (!load) {
return false;
}
args[i] = load;
}
}
return true;
}
bool Resolver::ShouldMaterializeArgument(const type::Type* parameter_ty) const {
const auto* param_el_ty = type::Type::DeepestElementOf(parameter_ty);
return param_el_ty && !param_el_ty->Is<type::AbstractNumeric>();
}
bool Resolver::Convert(const constant::Value*& c,
const type::Type* target_ty,
const Source& source) {
auto r = const_eval_.Convert(target_ty, c, source);
if (!r) {
return false;
}
c = r.Get();
return true;
}
template <size_t N>
utils::Result<utils::Vector<const constant::Value*, N>> Resolver::ConvertArguments(
const utils::Vector<const sem::Expression*, N>& args,
const sem::CallTarget* target) {
auto const_args = utils::Transform(args, [](auto* arg) { return arg->ConstantValue(); });
for (size_t i = 0, n = std::min(args.Length(), target->Parameters().Length()); i < n; i++) {
if (!Convert(const_args[i], target->Parameters()[i]->Type(),
args[i]->Declaration()->source)) {
return utils::Failure;
}
}
return const_args;
}
sem::Expression* Resolver::IndexAccessor(const ast::IndexAccessorExpression* expr) {
auto* idx = Load(Materialize(sem_.Get(expr->index)));
if (!idx) {
return nullptr;
}
const auto* obj = sem_.Get(expr->object);
if (idx->Stage() != sem::EvaluationStage::kConstant) {
// If the index is non-constant, then the resulting expression is non-constant, so we'll
// have to materialize the object. For example, consider:
// vec2(1, 2)[runtime-index]
obj = Materialize(obj);
}
if (!obj) {
return nullptr;
}
auto* obj_raw_ty = obj->Type();
auto* obj_ty = obj_raw_ty->UnwrapRef();
auto* ty = Switch(
obj_ty, //
[&](const type::Array* arr) { return arr->ElemType(); },
[&](const type::Vector* vec) { return vec->type(); },
[&](const type::Matrix* mat) {
return builder_->create<type::Vector>(mat->type(), mat->rows());
},
[&](Default) {
AddError("cannot index type '" + sem_.TypeNameOf(obj_ty) + "'", expr->source);
return nullptr;
});
if (ty == nullptr) {
return nullptr;
}
auto* idx_ty = idx->Type()->UnwrapRef();
if (!idx_ty->IsAnyOf<type::I32, type::U32>()) {
AddError("index must be of type 'i32' or 'u32', found: '" + sem_.TypeNameOf(idx_ty) + "'",
idx->Declaration()->source);
return nullptr;
}
// If we're extracting from a reference, we return a reference.
if (auto* ref = obj_raw_ty->As<type::Reference>()) {
ty = builder_->create<type::Reference>(ty, ref->AddressSpace(), ref->Access());
}
const constant::Value* val = nullptr;
auto stage = sem::EarliestStage(obj->Stage(), idx->Stage());
if (stage == sem::EvaluationStage::kConstant && skip_const_eval_.Contains(expr)) {
stage = sem::EvaluationStage::kNotEvaluated;
} else {
if (auto r = const_eval_.Index(obj, idx)) {
val = r.Get();
} else {
return nullptr;
}
}
bool has_side_effects = idx->HasSideEffects() || obj->HasSideEffects();
auto* sem = builder_->create<sem::IndexAccessorExpression>(
expr, ty, stage, obj, idx, current_statement_, std::move(val), has_side_effects,
obj->RootIdentifier());
sem->Behaviors() = idx->Behaviors() + obj->Behaviors();
return sem;
}
sem::Expression* Resolver::Bitcast(const ast::BitcastExpression* expr) {
auto* inner = Load(Materialize(sem_.Get(expr->expr)));
if (!inner) {
return nullptr;
}
auto* ty = Type(expr->type);
if (!ty) {
return nullptr;
}
if (!validator_.Bitcast(expr, ty)) {
return nullptr;
}
auto stage = inner->Stage();
if (stage == sem::EvaluationStage::kConstant && skip_const_eval_.Contains(expr)) {
stage = sem::EvaluationStage::kNotEvaluated;
}
const constant::Value* value = nullptr;
if (stage == sem::EvaluationStage::kConstant) {
if (auto r = const_eval_.Bitcast(ty, inner->ConstantValue(), expr->source)) {
value = r.Get();
} else {
return nullptr;
}
}
auto* sem = builder_->create<sem::Expression>(expr, ty, stage, current_statement_,
std::move(value), inner->HasSideEffects());
sem->Behaviors() = inner->Behaviors();
return sem;
}
sem::Call* Resolver::Call(const ast::CallExpression* expr) {
// A CallExpression can resolve to one of:
// * A function call.
// * A builtin call.
// * A type initializer.
// * A type conversion.
// Resolve all of the arguments, their types and the set of behaviors.
utils::Vector<const sem::Expression*, 8> args;
args.Reserve(expr->args.Length());
auto args_stage = sem::EvaluationStage::kConstant;
sem::Behaviors arg_behaviors;
for (size_t i = 0; i < expr->args.Length(); i++) {
auto* arg = sem_.Get(expr->args[i]);
if (!arg) {
return nullptr;
}
args.Push(arg);
args_stage = sem::EarliestStage(args_stage, arg->Stage());
arg_behaviors.Add(arg->Behaviors());
}
arg_behaviors.Remove(sem::Behavior::kNext);
// Did any arguments have side effects?
bool has_side_effects =
std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
// ct_init_or_conv is a helper for building either a sem::TypeInitializer or
// sem::TypeConversion call for a InitConvIntrinsic with an optional template argument type.
auto ct_init_or_conv = [&](InitConvIntrinsic ty, const type::Type* template_arg) -> sem::Call* {
auto arg_tys = utils::Transform(args, [](auto* arg) { return arg->Type(); });
auto ctor_or_conv =
intrinsic_table_->Lookup(ty, template_arg, arg_tys, args_stage, expr->source);
if (!ctor_or_conv.target) {
return nullptr;
}
if (!MaybeMaterializeAndLoadArguments(args, ctor_or_conv.target)) {
return nullptr;
}
const constant::Value* value = nullptr;
auto stage = sem::EarliestStage(ctor_or_conv.target->Stage(), args_stage);
if (stage == sem::EvaluationStage::kConstant && skip_const_eval_.Contains(expr)) {
stage = sem::EvaluationStage::kNotEvaluated;
}
if (stage == sem::EvaluationStage::kConstant) {
auto const_args = ConvertArguments(args, ctor_or_conv.target);
if (!const_args) {
return nullptr;
}
if (auto r = (const_eval_.*ctor_or_conv.const_eval_fn)(
ctor_or_conv.target->ReturnType(), const_args.Get(), expr->source)) {
value = r.Get();
} else {
return nullptr;
}
}
return builder_->create<sem::Call>(expr, ctor_or_conv.target, stage, std::move(args),
current_statement_, value, has_side_effects);
};
// arr_or_str_init is a helper for building a sem::TypeInitializer for an array or structure
// initializer call target.
auto arr_or_str_init = [&](const type::Type* ty,
const sem::CallTarget* call_target) -> sem::Call* {
if (!MaybeMaterializeAndLoadArguments(args, call_target)) {
return nullptr;
}
auto stage = args_stage; // The evaluation stage of the call
const constant::Value* value = nullptr; // The constant value for the call
if (stage == sem::EvaluationStage::kConstant) {
if (auto r = const_eval_.ArrayOrStructInit(ty, args)) {
value = r.Get();
} else {
return nullptr;
}
if (!value) {
// Constant evaluation failed.
// Can happen for expressions that will fail validation (later).
// Use the kRuntime EvaluationStage, as kConstant will trigger an assertion in
// the sem::Expression initializer, which checks that kConstant is paired with a
// constant value.
stage = sem::EvaluationStage::kRuntime;
}
}
return builder_->create<sem::Call>(expr, call_target, stage, std::move(args),
current_statement_, value, has_side_effects);
};
// ty_init_or_conv is a helper for building either a sem::TypeInitializer or
// sem::TypeConversion call for the given semantic type.
auto ty_init_or_conv = [&](const type::Type* ty) {
return Switch(
ty, //
[&](const type::Vector* v) {
return ct_init_or_conv(VectorInitConvIntrinsic(v->Width()), v->type());
},
[&](const type::Matrix* m) {
return ct_init_or_conv(MatrixInitConvIntrinsic(m->columns(), m->rows()), m->type());
},
[&](const type::I32*) { return ct_init_or_conv(InitConvIntrinsic::kI32, nullptr); },
[&](const type::U32*) { return ct_init_or_conv(InitConvIntrinsic::kU32, nullptr); },
[&](const type::F16*) {
return validator_.CheckF16Enabled(expr->source)
? ct_init_or_conv(InitConvIntrinsic::kF16, nullptr)
: nullptr;
},
[&](const type::F32*) { return ct_init_or_conv(InitConvIntrinsic::kF32, nullptr); },
[&](const type::Bool*) { return ct_init_or_conv(InitConvIntrinsic::kBool, nullptr); },
[&](const type::Array* arr) -> sem::Call* {
auto* call_target = array_inits_.GetOrCreate(
ArrayInitializerSig{{arr, args.Length(), args_stage}},
[&]() -> sem::TypeInitializer* {
auto params = utils::Transform(args, [&](auto, size_t i) {
return builder_->create<sem::Parameter>(
nullptr, // declaration
static_cast<uint32_t>(i), // index
arr->ElemType(), // type
ast::AddressSpace::kNone, // address_space
ast::Access::kUndefined);
});
return builder_->create<sem::TypeInitializer>(arr, std::move(params),
args_stage);
});
auto* call = arr_or_str_init(arr, call_target);
if (!call) {
return nullptr;
}
// Validation must occur after argument materialization in arr_or_str_init().
if (!validator_.ArrayInitializer(expr, arr)) {
return nullptr;
}
return call;
},
[&](const sem::Struct* str) -> sem::Call* {
auto* call_target = struct_inits_.GetOrCreate(
StructInitializerSig{{str, args.Length(), args_stage}},
[&]() -> sem::TypeInitializer* {
utils::Vector<const sem::Parameter*, 8> params;
params.Resize(std::min(args.Length(), str->Members().Length()));
for (size_t i = 0, n = params.Length(); i < n; i++) {
params[i] = builder_->create<sem::Parameter>(
nullptr, // declaration
static_cast<uint32_t>(i), // index
str->Members()[i]->Type(), // type
ast::AddressSpace::kNone, // address_space
ast::Access::kUndefined); // access
}
return builder_->create<sem::TypeInitializer>(str, std::move(params),
args_stage);
});
auto* call = arr_or_str_init(str, call_target);
if (!call) {
return nullptr;
}
// Validation must occur after argument materialization in arr_or_str_init().
if (!validator_.StructureInitializer(expr, str)) {
return nullptr;
}
return call;
},
[&](Default) {
AddError("type is not constructible", expr->source);
return nullptr;
});
};
// ast::CallExpression has a target which is either an ast::Type or an
// ast::IdentifierExpression
sem::Call* call = nullptr;
if (expr->target.type) {
// ast::CallExpression has an ast::Type as the target.
// This call is either a type initializer or type conversion.
call = Switch(
expr->target.type,
[&](const ast::Vector* v) -> sem::Call* {
Mark(v);
// vector element type must be inferred if it was not specified.
type::Type* template_arg = nullptr;
if (v->type) {
template_arg = Type(v->type);
if (!template_arg) {
return nullptr;
}
}
if (auto* c = ct_init_or_conv(VectorInitConvIntrinsic(v->width), template_arg)) {
builder_->Sem().Add(expr->target.type, c->Target()->ReturnType());
return c;
}
return nullptr;
},
[&](const ast::Matrix* m) -> sem::Call* {
Mark(m);
// matrix element type must be inferred if it was not specified.
type::Type* template_arg = nullptr;
if (m->type) {
template_arg = Type(m->type);
if (!template_arg) {
return nullptr;
}
}
if (auto* c = ct_init_or_conv(MatrixInitConvIntrinsic(m->columns, m->rows),
template_arg)) {
builder_->Sem().Add(expr->target.type, c->Target()->ReturnType());
return c;
}
return nullptr;
},
[&](const ast::Array* a) -> sem::Call* {
Mark(a);
// array element type must be inferred if it was not specified.
const type::ArrayCount* el_count = nullptr;
const type::Type* el_ty = nullptr;
if (a->type) {
el_ty = Type(a->type);
if (!el_ty) {
return nullptr;
}
if (!a->count) {
AddError("cannot construct a runtime-sized array", expr->source);
return nullptr;
}
el_count = ArrayCount(a->count);
if (!el_count) {
return nullptr;
}
// Note: validation later will detect any mismatches between explicit array
// size and number of initializer expressions.
} else {
el_count = builder_->create<type::ConstantArrayCount>(
static_cast<uint32_t>(args.Length()));
auto arg_tys =
utils::Transform(args, [](auto* arg) { return arg->Type()->UnwrapRef(); });
el_ty = type::Type::Common(arg_tys);
if (!el_ty) {
AddError(
"cannot infer common array element type from initializer arguments",
expr->source);
utils::Hashset<const type::Type*, 8> types;
for (size_t i = 0; i < args.Length(); i++) {
if (types.Add(args[i]->Type())) {
AddNote("argument " + std::to_string(i) + " is of type '" +
sem_.TypeNameOf(args[i]->Type()) + "'",
args[i]->Declaration()->source);
}
}
return nullptr;
}
}
uint32_t explicit_stride = 0;
if (!ArrayAttributes(a->attributes, el_ty, explicit_stride)) {
return nullptr;
}
auto* arr = Array(a->type ? a->type->source : a->source,
a->count ? a->count->source : a->source, //
el_ty, el_count, explicit_stride);
if (!arr) {
return nullptr;
}
builder_->Sem().Add(a, arr);
return ty_init_or_conv(arr);
},
[&](const ast::Type* ast) -> sem::Call* {
// Handler for AST types that do not have an optional element type.
if (auto* ty = Type(ast)) {
return ty_init_or_conv(ty);
}
return nullptr;
},
[&](Default) {
TINT_ICE(Resolver, diagnostics_)
<< expr->source << " unhandled CallExpression target:\n"
<< "type: "
<< (expr->target.type ? expr->target.type->TypeInfo().name : "<null>");
return nullptr;
});
} else {
// ast::CallExpression has an ast::IdentifierExpression as the target.
// This call is either a function call, builtin call, type initializer or type
// conversion.
auto* ident = expr->target.name;
Mark(ident);
if (auto* resolved = sem_.ResolvedSymbol<type::Type>(ident)) {
// A type initializer or conversions.
// Note: Unlike the code path where we're resolving the call target from an
// ast::Type, all types must already have the element type explicitly specified,
// so there's no need to infer element types.
return ty_init_or_conv(resolved);
}
auto* resolved = sem_.ResolvedSymbol<sem::Node>(ident);
call = Switch<sem::Call*>(
resolved, //
[&](sem::Function* func) { return FunctionCall(expr, func, args, arg_behaviors); },
[&](sem::Variable* var) {
auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
AddError("cannot call variable '" + name + "'", ident->source);
AddNote("'" + name + "' declared here", var->Declaration()->source);
return nullptr;
},
[&](Default) -> sem::Call* {
auto name = builder_->Symbols().NameFor(ident->symbol);
if (auto builtin_type = sem::ParseBuiltinType(name);
builtin_type != sem::BuiltinType::kNone) {
return BuiltinCall(expr, builtin_type, args);
}
if (auto* alias = ShortName(ident->symbol, ident->source)) {
return ty_init_or_conv(alias);
}
return nullptr;
});
}
if (!call) {
return nullptr;
}
return validator_.Call(call, current_statement_) ? call : nullptr;
}
template <size_t N>
sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
sem::BuiltinType builtin_type,
utils::Vector<const sem::Expression*, N>& args) {
auto arg_stage = sem::EvaluationStage::kConstant;
for (auto* arg : args) {
arg_stage = sem::EarliestStage(arg_stage, arg->Stage());
}
IntrinsicTable::Builtin builtin;
{
auto arg_tys = utils::Transform(args, [](auto* arg) { return arg->Type(); });
builtin = intrinsic_table_->Lookup(builtin_type, arg_tys, arg_stage, expr->source);
if (!builtin.sem) {
return nullptr;
}
}
if (builtin_type == sem::BuiltinType::kTintMaterialize) {
args[0] = Materialize(args[0]);
if (!args[0]) {
return nullptr;
}
} else {
// Materialize arguments if the parameter type is not abstract
if (!MaybeMaterializeAndLoadArguments(args, builtin.sem)) {
return nullptr;
}
}
if (builtin.sem->IsDeprecated()) {
AddWarning("use of deprecated builtin", expr->source);
}
// If the builtin is @const, and all arguments have constant values, evaluate the builtin
// now.
const constant::Value* value = nullptr;
auto stage = sem::EarliestStage(arg_stage, builtin.sem->Stage());
if (stage == sem::EvaluationStage::kConstant && skip_const_eval_.Contains(expr)) {
stage = sem::EvaluationStage::kNotEvaluated;
}
if (stage == sem::EvaluationStage::kConstant) {
auto const_args = ConvertArguments(args, builtin.sem);
if (!const_args) {
return nullptr;
}
if (auto r = (const_eval_.*builtin.const_eval_fn)(builtin.sem->ReturnType(),
const_args.Get(), expr->source)) {
value = r.Get();
} else {
return nullptr;
}
}
bool has_side_effects =
builtin.sem->HasSideEffects() ||
std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
auto* call = builder_->create<sem::Call>(expr, builtin.sem, stage, std::move(args),
current_statement_, value, has_side_effects);
if (current_function_) {
current_function_->AddDirectlyCalledBuiltin(builtin.sem);
current_function_->AddDirectCall(call);
}
if (!validator_.RequiredExtensionForBuiltinFunction(call)) {
return nullptr;
}
if (IsTextureBuiltin(builtin_type)) {
if (!validator_.TextureBuiltinFunction(call)) {
return nullptr;
}
CollectTextureSamplerPairs(builtin.sem, call->Arguments());
}
if (!validator_.BuiltinCall(call)) {
return nullptr;
}
return call;
}
type::Type* Resolver::ShortName(Symbol sym, const Source& source) const {
auto name = builder_->Symbols().NameFor(sym);
auto& b = *builder_;
auto vec_f32 = [&](uint32_t n) { return b.create<type::Vector>(b.create<type::F32>(), n); };
auto vec_f16 = [&](uint32_t n) { return b.create<type::Vector>(b.create<type::F16>(), n); };
switch (type::ParseShortName(name)) {
case type::ShortName::kMat2X2F:
return b.create<type::Matrix>(vec_f32(2u), 2u);
case type::ShortName::kMat2X3F:
return b.create<type::Matrix>(vec_f32(3u), 2u);
case type::ShortName::kMat2X4F:
return b.create<type::Matrix>(vec_f32(4u), 2u);
case type::ShortName::kMat3X2F:
return b.create<type::Matrix>(vec_f32(2u), 3u);
case type::ShortName::kMat3X3F:
return b.create<type::Matrix>(vec_f32(3u), 3u);
case type::ShortName::kMat3X4F:
return b.create<type::Matrix>(vec_f32(4u), 3u);
case type::ShortName::kMat4X2F:
return b.create<type::Matrix>(vec_f32(2u), 4u);
case type::ShortName::kMat4X3F:
return b.create<type::Matrix>(vec_f32(3u), 4u);
case type::ShortName::kMat4X4F:
return b.create<type::Matrix>(vec_f32(4u), 4u);
case type::ShortName::kMat2X2H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(2u), 2u)
: nullptr;
case type::ShortName::kMat2X3H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(3u), 2u)
: nullptr;
case type::ShortName::kMat2X4H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(4u), 2u)
: nullptr;
case type::ShortName::kMat3X2H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(2u), 3u)
: nullptr;
case type::ShortName::kMat3X3H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(3u), 3u)
: nullptr;
case type::ShortName::kMat3X4H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(4u), 3u)
: nullptr;
case type::ShortName::kMat4X2H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(2u), 4u)
: nullptr;
case type::ShortName::kMat4X3H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(3u), 4u)
: nullptr;
case type::ShortName::kMat4X4H:
return validator_.CheckF16Enabled(source) ? b.create<type::Matrix>(vec_f16(4u), 4u)
: nullptr;
case type::ShortName::kVec2F:
return vec_f32(2u);
case type::ShortName::kVec3F:
return vec_f32(3u);
case type::ShortName::kVec4F:
return vec_f32(4u);
case type::ShortName::kVec2H:
return validator_.CheckF16Enabled(source) ? vec_f16(2u) : nullptr;
case type::ShortName::kVec3H:
return validator_.CheckF16Enabled(source) ? vec_f16(3u) : nullptr;
case type::ShortName::kVec4H:
return validator_.CheckF16Enabled(source) ? vec_f16(4u) : nullptr;
case type::ShortName::kVec2I:
return b.create<type::Vector>(b.create<type::I32>(), 2u);
case type::ShortName::kVec3I:
return b.create<type::Vector>(b.create<type::I32>(), 3u);
case type::ShortName::kVec4I:
return b.create<type::Vector>(b.create<type::I32>(), 4u);
case type::ShortName::kVec2U:
return b.create<type::Vector>(b.create<type::U32>(), 2u);
case type::ShortName::kVec3U:
return b.create<type::Vector>(b.create<type::U32>(), 3u);
case type::ShortName::kVec4U:
return b.create<type::Vector>(b.create<type::U32>(), 4u);
case type::ShortName::kUndefined:
break;
}
TINT_ICE(Resolver, diagnostics_) << source << " unhandled type short name '" << name << "'";
return nullptr;
}
void Resolver::CollectTextureSamplerPairs(const sem::Builtin* builtin,
utils::VectorRef<const sem::Expression*> args) const {
// Collect a texture/sampler pair for this builtin.
const auto& signature = builtin->Signature();
int texture_index = signature.IndexOf(sem::ParameterUsage::kTexture);
if (texture_index == -1) {
TINT_ICE(Resolver, diagnostics_) << "texture builtin without texture parameter";
}
if (auto* user =
args[static_cast<size_t>(texture_index)]->UnwrapLoad()->As<sem::VariableUser>()) {
auto* texture = user->Variable();
if (!texture->Type()->UnwrapRef()->Is<type::StorageTexture>()) {
int sampler_index = signature.IndexOf(sem::ParameterUsage::kSampler);
const sem::Variable* sampler = sampler_index != -1
? args[static_cast<size_t>(sampler_index)]
->UnwrapLoad()
->As<sem::VariableUser>()
->Variable()
: nullptr;
current_function_->AddTextureSamplerPair(texture, sampler);
}
}
}
template <size_t N>
sem::Call* Resolver::FunctionCall(const ast::CallExpression* expr,
sem::Function* target,
utils::Vector<const sem::Expression*, N>& args,
sem::Behaviors arg_behaviors) {
auto sym = expr->target.name->symbol;
auto name = builder_->Symbols().NameFor(sym);
if (!MaybeMaterializeAndLoadArguments(args, target)) {
return nullptr;
}
// TODO(crbug.com/tint/1420): For now, assume all function calls have side
// effects.
bool has_side_effects = true;
auto* call = builder_->create<sem::Call>(expr, target, sem::EvaluationStage::kRuntime,
std::move(args), current_statement_,
/* constant_value */ nullptr, has_side_effects);
target->AddCallSite(call);
call->Behaviors() = arg_behaviors + target->Behaviors();
if (!validator_.FunctionCall(call, current_statement_)) {
return nullptr;
}
if (current_function_) {
// Note: Requires called functions to be resolved first.
// This is currently guaranteed as functions must be declared before
// use.
current_function_->AddTransitivelyCalledFunction(target);
current_function_->AddDirectCall(call);
for (auto* transitive_call : target->TransitivelyCalledFunctions()) {