| // 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/resolver/resolver.h" |
| |
| #include <algorithm> |
| #include <cmath> |
| #include <iomanip> |
| #include <limits> |
| #include <utility> |
| |
| #include "src/ast/alias.h" |
| #include "src/ast/array.h" |
| #include "src/ast/assignment_statement.h" |
| #include "src/ast/bitcast_expression.h" |
| #include "src/ast/break_statement.h" |
| #include "src/ast/call_statement.h" |
| #include "src/ast/continue_statement.h" |
| #include "src/ast/depth_texture.h" |
| #include "src/ast/disable_validation_decoration.h" |
| #include "src/ast/discard_statement.h" |
| #include "src/ast/fallthrough_statement.h" |
| #include "src/ast/for_loop_statement.h" |
| #include "src/ast/if_statement.h" |
| #include "src/ast/internal_decoration.h" |
| #include "src/ast/interpolate_decoration.h" |
| #include "src/ast/loop_statement.h" |
| #include "src/ast/matrix.h" |
| #include "src/ast/override_decoration.h" |
| #include "src/ast/pointer.h" |
| #include "src/ast/return_statement.h" |
| #include "src/ast/sampled_texture.h" |
| #include "src/ast/sampler.h" |
| #include "src/ast/storage_texture.h" |
| #include "src/ast/struct_block_decoration.h" |
| #include "src/ast/switch_statement.h" |
| #include "src/ast/traverse_expressions.h" |
| #include "src/ast/type_name.h" |
| #include "src/ast/unary_op_expression.h" |
| #include "src/ast/variable_decl_statement.h" |
| #include "src/ast/vector.h" |
| #include "src/ast/workgroup_decoration.h" |
| #include "src/sem/array.h" |
| #include "src/sem/atomic_type.h" |
| #include "src/sem/call.h" |
| #include "src/sem/depth_multisampled_texture_type.h" |
| #include "src/sem/depth_texture_type.h" |
| #include "src/sem/for_loop_statement.h" |
| #include "src/sem/function.h" |
| #include "src/sem/if_statement.h" |
| #include "src/sem/loop_statement.h" |
| #include "src/sem/member_accessor_expression.h" |
| #include "src/sem/multisampled_texture_type.h" |
| #include "src/sem/pointer_type.h" |
| #include "src/sem/reference_type.h" |
| #include "src/sem/sampled_texture_type.h" |
| #include "src/sem/sampler_type.h" |
| #include "src/sem/statement.h" |
| #include "src/sem/storage_texture_type.h" |
| #include "src/sem/struct.h" |
| #include "src/sem/switch_statement.h" |
| #include "src/sem/variable.h" |
| #include "src/utils/defer.h" |
| #include "src/utils/get_or_create.h" |
| #include "src/utils/math.h" |
| #include "src/utils/reverse.h" |
| #include "src/utils/scoped_assignment.h" |
| |
| namespace tint { |
| namespace resolver { |
| namespace { |
| |
| using IntrinsicType = tint::sem::IntrinsicType; |
| |
| bool IsValidStorageTextureDimension(ast::TextureDimension dim) { |
| switch (dim) { |
| case ast::TextureDimension::k1d: |
| case ast::TextureDimension::k2d: |
| case ast::TextureDimension::k2dArray: |
| case ast::TextureDimension::k3d: |
| return true; |
| default: |
| return false; |
| } |
| } |
| |
| bool IsValidStorageTextureImageFormat(ast::ImageFormat format) { |
| switch (format) { |
| case ast::ImageFormat::kR32Uint: |
| case ast::ImageFormat::kR32Sint: |
| case ast::ImageFormat::kR32Float: |
| case ast::ImageFormat::kRg32Uint: |
| case ast::ImageFormat::kRg32Sint: |
| case ast::ImageFormat::kRg32Float: |
| case ast::ImageFormat::kRgba8Unorm: |
| case ast::ImageFormat::kRgba8Snorm: |
| case ast::ImageFormat::kRgba8Uint: |
| case ast::ImageFormat::kRgba8Sint: |
| case ast::ImageFormat::kRgba16Uint: |
| case ast::ImageFormat::kRgba16Sint: |
| case ast::ImageFormat::kRgba16Float: |
| case ast::ImageFormat::kRgba32Uint: |
| case ast::ImageFormat::kRgba32Sint: |
| case ast::ImageFormat::kRgba32Float: |
| return true; |
| default: |
| return false; |
| } |
| } |
| |
| /// @returns true if the decoration list contains a |
| /// ast::DisableValidationDecoration with the validation mode equal to |
| /// `validation` |
| bool IsValidationDisabled(const ast::DecorationList& decorations, |
| ast::DisabledValidation validation) { |
| for (auto* decoration : decorations) { |
| if (auto* dv = decoration->As<ast::DisableValidationDecoration>()) { |
| if (dv->validation == validation) { |
| return true; |
| } |
| } |
| } |
| return false; |
| } |
| |
| /// @returns true if the decoration list does not contains a |
| /// ast::DisableValidationDecoration with the validation mode equal to |
| /// `validation` |
| bool IsValidationEnabled(const ast::DecorationList& decorations, |
| ast::DisabledValidation validation) { |
| return !IsValidationDisabled(decorations, validation); |
| } |
| |
| // Helper to stringify a pipeline IO decoration. |
| std::string deco_to_str(const ast::Decoration* deco) { |
| std::stringstream str; |
| if (auto* builtin = deco->As<ast::BuiltinDecoration>()) { |
| str << "builtin(" << builtin->builtin << ")"; |
| } else if (auto* location = deco->As<ast::LocationDecoration>()) { |
| str << "location(" << location->value << ")"; |
| } |
| return str.str(); |
| } |
| } // namespace |
| |
| Resolver::Resolver(ProgramBuilder* builder) |
| : builder_(builder), |
| diagnostics_(builder->Diagnostics()), |
| intrinsic_table_(IntrinsicTable::Create(*builder)) {} |
| |
| Resolver::~Resolver() = default; |
| |
| void Resolver::set_referenced_from_function_if_needed(VariableInfo* var, |
| bool local) { |
| if (current_function_ == nullptr) { |
| return; |
| } |
| |
| if (var->kind != VariableKind::kGlobal) { |
| return; |
| } |
| |
| current_function_->referenced_module_vars.add(var); |
| if (local) { |
| current_function_->local_referenced_module_vars.add(var); |
| } |
| } |
| |
| bool Resolver::Resolve() { |
| if (builder_->Diagnostics().contains_errors()) { |
| return false; |
| } |
| |
| bool result = ResolveInternal(); |
| |
| if (!result && !diagnostics_.contains_errors()) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "resolving failed, but no error was raised"; |
| return false; |
| } |
| |
| // Even if resolving failed, create all the semantic nodes for information we |
| // did generate. |
| CreateSemanticNodes(); |
| |
| return result; |
| } |
| |
| // https://gpuweb.github.io/gpuweb/wgsl/#plain-types-section |
| bool Resolver::IsPlain(const sem::Type* type) const { |
| return type->is_scalar() || type->Is<sem::Atomic>() || |
| type->Is<sem::Vector>() || type->Is<sem::Matrix>() || |
| type->Is<sem::Array>() || type->Is<sem::Struct>(); |
| } |
| |
| // https://gpuweb.github.io/gpuweb/wgsl.html#storable-types |
| bool Resolver::IsStorable(const sem::Type* type) const { |
| return IsPlain(type) || type->Is<sem::Texture>() || type->Is<sem::Sampler>(); |
| } |
| |
| // https://gpuweb.github.io/gpuweb/wgsl.html#host-shareable-types |
| bool Resolver::IsHostShareable(const sem::Type* type) const { |
| if (type->IsAnyOf<sem::I32, sem::U32, sem::F32>()) { |
| return true; |
| } |
| if (auto* vec = type->As<sem::Vector>()) { |
| return IsHostShareable(vec->type()); |
| } |
| if (auto* mat = type->As<sem::Matrix>()) { |
| return IsHostShareable(mat->type()); |
| } |
| if (auto* arr = type->As<sem::Array>()) { |
| return IsHostShareable(arr->ElemType()); |
| } |
| if (auto* str = type->As<sem::Struct>()) { |
| for (auto* member : str->Members()) { |
| if (!IsHostShareable(member->Type())) { |
| return false; |
| } |
| } |
| return true; |
| } |
| if (auto* atomic = type->As<sem::Atomic>()) { |
| return IsHostShareable(atomic->Type()); |
| } |
| return false; |
| } |
| |
| bool Resolver::ResolveInternal() { |
| Mark(&builder_->AST()); |
| |
| // Process everything else in the order they appear in the module. This is |
| // necessary for validation of use-before-declaration. |
| for (auto* decl : builder_->AST().GlobalDeclarations()) { |
| if (auto* td = decl->As<ast::TypeDecl>()) { |
| Mark(td); |
| if (!TypeDecl(td)) { |
| return false; |
| } |
| } else if (auto* func = decl->As<ast::Function>()) { |
| Mark(func); |
| if (!Function(func)) { |
| return false; |
| } |
| } else if (auto* var = decl->As<ast::Variable>()) { |
| Mark(var); |
| if (!GlobalVariable(var)) { |
| return false; |
| } |
| } else { |
| TINT_UNREACHABLE(Resolver, diagnostics_) |
| << "unhandled global declaration: " << decl->TypeInfo().name; |
| return false; |
| } |
| } |
| |
| AllocateOverridableConstantIds(); |
| |
| if (!ValidatePipelineStages()) { |
| return false; |
| } |
| |
| bool result = true; |
| for (auto* node : builder_->ASTNodes().Objects()) { |
| if (marked_.count(node) == 0) { |
| 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; |
| } |
| |
| sem::Type* Resolver::Type(const ast::Type* ty) { |
| Mark(ty); |
| auto* s = [&]() -> sem::Type* { |
| if (ty->Is<ast::Void>()) { |
| return builder_->create<sem::Void>(); |
| } |
| if (ty->Is<ast::Bool>()) { |
| return builder_->create<sem::Bool>(); |
| } |
| if (ty->Is<ast::I32>()) { |
| return builder_->create<sem::I32>(); |
| } |
| if (ty->Is<ast::U32>()) { |
| return builder_->create<sem::U32>(); |
| } |
| if (ty->Is<ast::F32>()) { |
| return builder_->create<sem::F32>(); |
| } |
| if (auto* t = ty->As<ast::Vector>()) { |
| if (auto* el = Type(t->type)) { |
| if (auto* vector = builder_->create<sem::Vector>(el, t->width)) { |
| if (ValidateVector(vector, t->source)) { |
| return vector; |
| } |
| } |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::Matrix>()) { |
| if (auto* el = Type(t->type)) { |
| if (auto* column_type = builder_->create<sem::Vector>(el, t->rows)) { |
| if (auto* matrix = |
| builder_->create<sem::Matrix>(column_type, t->columns)) { |
| if (ValidateMatrix(matrix, t->source)) { |
| return matrix; |
| } |
| } |
| } |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::Array>()) { |
| return Array(t); |
| } |
| if (auto* t = ty->As<ast::Atomic>()) { |
| if (auto* el = Type(t->type)) { |
| auto* a = builder_->create<sem::Atomic>(el); |
| if (!ValidateAtomic(t, a)) { |
| return nullptr; |
| } |
| return a; |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::Pointer>()) { |
| if (auto* el = Type(t->type)) { |
| auto access = t->access; |
| if (access == ast::kUndefined) { |
| access = DefaultAccessForStorageClass(t->storage_class); |
| } |
| return builder_->create<sem::Pointer>(el, t->storage_class, access); |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::Sampler>()) { |
| return builder_->create<sem::Sampler>(t->kind); |
| } |
| if (auto* t = ty->As<ast::SampledTexture>()) { |
| if (auto* el = Type(t->type)) { |
| return builder_->create<sem::SampledTexture>(t->dim, el); |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::MultisampledTexture>()) { |
| if (auto* el = Type(t->type)) { |
| return builder_->create<sem::MultisampledTexture>(t->dim, el); |
| } |
| return nullptr; |
| } |
| if (auto* t = ty->As<ast::DepthTexture>()) { |
| return builder_->create<sem::DepthTexture>(t->dim); |
| } |
| if (auto* t = ty->As<ast::DepthMultisampledTexture>()) { |
| return builder_->create<sem::DepthMultisampledTexture>(t->dim); |
| } |
| if (auto* t = ty->As<ast::StorageTexture>()) { |
| if (auto* el = Type(t->type)) { |
| if (!ValidateStorageTexture(t)) { |
| return nullptr; |
| } |
| return builder_->create<sem::StorageTexture>(t->dim, t->format, |
| t->access, el); |
| } |
| return nullptr; |
| } |
| if (ty->As<ast::ExternalTexture>()) { |
| return builder_->create<sem::ExternalTexture>(); |
| } |
| if (auto* t = ty->As<ast::TypeName>()) { |
| auto it = named_type_info_.find(t->name); |
| if (it == named_type_info_.end()) { |
| AddError("unknown type '" + builder_->Symbols().NameFor(t->name) + "'", |
| t->source); |
| return nullptr; |
| } |
| return it->second.sem; |
| } |
| TINT_UNREACHABLE(Resolver, diagnostics_) |
| << "Unhandled ast::Type: " << ty->TypeInfo().name; |
| return nullptr; |
| }(); |
| |
| if (s) { |
| builder_->Sem().Add(ty, s); |
| } |
| return s; |
| } |
| |
| bool Resolver::ValidateAtomic(const ast::Atomic* a, const sem::Atomic* s) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types |
| // T must be either u32 or i32. |
| if (!s->Type()->IsAnyOf<sem::U32, sem::I32>()) { |
| AddError("atomic only supports i32 or u32 types", |
| a->type ? a->type->source : a->source); |
| return false; |
| } |
| return true; |
| } |
| |
| bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) { |
| switch (t->access) { |
| case ast::Access::kWrite: |
| break; |
| case ast::Access::kUndefined: |
| AddError("storage texture missing access control", t->source); |
| return false; |
| default: |
| AddError("storage textures currently only support 'write' access control", |
| t->source); |
| return false; |
| } |
| |
| if (!IsValidStorageTextureDimension(t->dim)) { |
| AddError("cube dimensions for storage textures are not supported", |
| t->source); |
| return false; |
| } |
| |
| if (!IsValidStorageTextureImageFormat(t->format)) { |
| AddError( |
| "image format must be one of the texel formats specified for storage " |
| "textues in https://gpuweb.github.io/gpuweb/wgsl/#texel-formats", |
| t->source); |
| return false; |
| } |
| return true; |
| } |
| |
| Resolver::VariableInfo* Resolver::Variable(const ast::Variable* var, |
| VariableKind kind, |
| uint32_t index /* = 0 */) { |
| if (variable_to_info_.count(var)) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "Variable " << builder_->Symbols().NameFor(var->symbol) |
| << " already resolved"; |
| return nullptr; |
| } |
| |
| std::string type_name; |
| const sem::Type* storage_type = nullptr; |
| |
| // If the variable has a declared type, resolve it. |
| if (auto* ty = var->type) { |
| type_name = ty->FriendlyName(builder_->Symbols()); |
| storage_type = Type(ty); |
| if (!storage_type) { |
| return nullptr; |
| } |
| } |
| |
| std::string rhs_type_name; |
| const sem::Type* rhs_type = nullptr; |
| |
| // Does the variable have a constructor? |
| if (auto* ctor = var->constructor) { |
| if (!Expression(var->constructor)) { |
| return nullptr; |
| } |
| |
| // Fetch the constructor's type |
| rhs_type_name = TypeNameOf(ctor); |
| rhs_type = TypeOf(ctor); |
| if (!rhs_type) { |
| return nullptr; |
| } |
| |
| // If the variable has no declared type, infer it from the RHS |
| if (!storage_type) { |
| if (!var->is_const && kind == VariableKind::kGlobal) { |
| AddError("global var declaration must specify a type", var->source); |
| return nullptr; |
| } |
| |
| type_name = rhs_type_name; |
| storage_type = rhs_type->UnwrapRef(); // Implicit load of RHS |
| } |
| } else if (var->is_const && kind != VariableKind::kParameter && |
| !ast::HasDecoration<ast::OverrideDecoration>(var->decorations)) { |
| AddError("let declaration must have an initializer", var->source); |
| return nullptr; |
| } else if (!var->type) { |
| AddError( |
| (kind == VariableKind::kGlobal) |
| ? "module scope var declaration requires a type and initializer" |
| : "function scope var declaration requires a type or initializer", |
| var->source); |
| return nullptr; |
| } |
| |
| if (!storage_type) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "failed to determine storage type for variable '" + |
| builder_->Symbols().NameFor(var->symbol) + "'\n" |
| << "Source: " << var->source; |
| return nullptr; |
| } |
| |
| auto storage_class = var->declared_storage_class; |
| if (storage_class == ast::StorageClass::kNone && !var->is_const) { |
| // No declared storage class. Infer from usage / type. |
| if (kind == VariableKind::kLocal) { |
| storage_class = ast::StorageClass::kFunction; |
| } else if (storage_type->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 storage class decoration. The |
| // storage class will always be handle. |
| storage_class = ast::StorageClass::kUniformConstant; |
| } |
| } |
| |
| auto access = var->declared_access; |
| if (access == ast::Access::kUndefined) { |
| access = DefaultAccessForStorageClass(storage_class); |
| } |
| |
| auto* type = storage_type; |
| if (!var->is_const) { |
| // Variable declaration. Unlike `let`, `var` has storage. |
| // Variables are always of a reference type to the declared storage type. |
| type = |
| builder_->create<sem::Reference>(storage_type, storage_class, access); |
| } |
| |
| if (rhs_type && |
| !ValidateVariableConstructor(var, storage_class, storage_type, type_name, |
| rhs_type, rhs_type_name)) { |
| return nullptr; |
| } |
| |
| auto* info = |
| variable_infos_.Create(var, const_cast<sem::Type*>(type), type_name, |
| storage_class, access, kind, index); |
| variable_to_info_.emplace(var, info); |
| |
| return info; |
| } |
| |
| ast::Access Resolver::DefaultAccessForStorageClass( |
| ast::StorageClass storage_class) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#storage-class |
| switch (storage_class) { |
| case ast::StorageClass::kStorage: |
| case ast::StorageClass::kUniform: |
| case ast::StorageClass::kUniformConstant: |
| return ast::Access::kRead; |
| default: |
| break; |
| } |
| return ast::Access::kReadWrite; |
| } |
| |
| void Resolver::AllocateOverridableConstantIds() { |
| // The next pipeline constant ID to try to allocate. |
| uint16_t next_constant_id = 0; |
| |
| // 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* var = decl->As<ast::Variable>(); |
| if (!var) { |
| continue; |
| } |
| auto* override_deco = |
| ast::GetDecoration<ast::OverrideDecoration>(var->decorations); |
| if (!override_deco) { |
| continue; |
| } |
| |
| uint16_t constant_id; |
| if (override_deco->has_value) { |
| constant_id = static_cast<uint16_t>(override_deco->value); |
| } else { |
| // No ID was specified, so allocate the next available ID. |
| constant_id = next_constant_id; |
| while (constant_ids_.count(constant_id)) { |
| if (constant_id == UINT16_MAX) { |
| TINT_ICE(Resolver, builder_->Diagnostics()) |
| << "no more pipeline constant IDs available"; |
| return; |
| } |
| constant_id++; |
| } |
| next_constant_id = constant_id + 1; |
| } |
| |
| variable_to_info_[var]->constant_id = constant_id; |
| } |
| } |
| |
| bool Resolver::ValidateVariableConstructor(const ast::Variable* var, |
| ast::StorageClass storage_class, |
| const sem::Type* storage_type, |
| const std::string& type_name, |
| const sem::Type* rhs_type, |
| const std::string& rhs_type_name) { |
| auto* value_type = rhs_type->UnwrapRef(); // Implicit load of RHS |
| |
| // Value type has to match storage type |
| if (storage_type != value_type) { |
| std::string decl = var->is_const ? "let" : "var"; |
| AddError("cannot initialize " + decl + " of type '" + type_name + |
| "' with value of type '" + rhs_type_name + "'", |
| var->source); |
| return false; |
| } |
| |
| if (!var->is_const) { |
| switch (storage_class) { |
| case ast::StorageClass::kPrivate: |
| case ast::StorageClass::kFunction: |
| break; // Allowed an initializer |
| default: |
| // https://gpuweb.github.io/gpuweb/wgsl/#var-and-let |
| // Optionally has an initializer expression, if the variable is in the |
| // private or function storage classes. |
| AddError("var of storage class '" + |
| std::string(ast::ToString(storage_class)) + |
| "' cannot have an initializer. var initializers are only " |
| "supported for the storage classes " |
| "'private' and 'function'", |
| var->source); |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::GlobalVariable(const ast::Variable* var) { |
| if (!ValidateNoDuplicateDefinition(var->symbol, var->source, |
| /* check_global_scope_only */ true)) { |
| return false; |
| } |
| |
| auto* info = Variable(var, VariableKind::kGlobal); |
| if (!info) { |
| return false; |
| } |
| variable_stack_.Set(var->symbol, info); |
| |
| if (!var->is_const && info->storage_class == ast::StorageClass::kNone) { |
| AddError("global variables must have a storage class", var->source); |
| return false; |
| } |
| if (var->is_const && !(info->storage_class == ast::StorageClass::kNone)) { |
| AddError("global constants shouldn't have a storage class", var->source); |
| return false; |
| } |
| |
| for (auto* deco : var->decorations) { |
| Mark(deco); |
| |
| if (auto* override_deco = deco->As<ast::OverrideDecoration>()) { |
| // Track the constant IDs that are specified in the shader. |
| if (override_deco->has_value) { |
| constant_ids_.emplace(override_deco->value, info); |
| } |
| } |
| } |
| |
| if (!ValidateNoDuplicateDecorations(var->decorations)) { |
| return false; |
| } |
| |
| if (auto bp = var->BindingPoint()) { |
| info->binding_point = {bp.group->value, bp.binding->value}; |
| } |
| |
| if (!ValidateGlobalVariable(info)) { |
| return false; |
| } |
| |
| if (!ApplyStorageClassUsageToType( |
| info->storage_class, const_cast<sem::Type*>(info->type->UnwrapRef()), |
| var->source)) { |
| AddNote("while instantiating variable " + |
| builder_->Symbols().NameFor(var->symbol), |
| var->source); |
| return false; |
| } |
| |
| // TODO(bclayton): Call this at the end of resolve on all uniform and storage |
| // referenced structs |
| if (!ValidateStorageClassLayout(info)) { |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateStorageClassLayout(const sem::Struct* str, |
| ast::StorageClass sc) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints |
| |
| auto is_uniform_struct_or_array = [sc](const sem::Type* ty) { |
| return sc == ast::StorageClass::kUniform && |
| ty->IsAnyOf<sem::Array, sem::Struct>(); |
| }; |
| |
| auto is_uniform_struct = [sc](const sem::Type* ty) { |
| return sc == ast::StorageClass::kUniform && ty->Is<sem::Struct>(); |
| }; |
| |
| auto required_alignment_of = [&](const sem::Type* ty) { |
| uint32_t actual_align = ty->Align(); |
| uint32_t required_align = actual_align; |
| if (is_uniform_struct_or_array(ty)) { |
| required_align = utils::RoundUp(16u, actual_align); |
| } |
| return required_align; |
| }; |
| |
| auto member_name_of = [this](const sem::StructMember* sm) { |
| return builder_->Symbols().NameFor(sm->Declaration()->symbol); |
| }; |
| |
| auto type_name_of = [this](const sem::StructMember* sm) { |
| return sm->Declaration()->type->FriendlyName(builder_->Symbols()); |
| }; |
| |
| // TODO(amaiorano): Output struct and member decorations so that this output |
| // can be copied verbatim back into source |
| auto get_struct_layout_string = [&](const sem::Struct* st) -> std::string { |
| std::stringstream ss; |
| |
| if (st->Members().empty()) { |
| TINT_ICE(Resolver, diagnostics_) << "Validation should have ensured that " |
| "structs have at least one member"; |
| return {}; |
| } |
| const auto* const last_member = st->Members().back(); |
| const uint32_t last_member_struct_padding_offset = |
| last_member->Offset() + last_member->Size(); |
| |
| // Compute max widths to align output |
| const auto offset_w = |
| static_cast<int>(::log10(last_member_struct_padding_offset)) + 1; |
| const auto size_w = static_cast<int>(::log10(st->Size())) + 1; |
| const auto align_w = static_cast<int>(::log10(st->Align())) + 1; |
| |
| auto print_struct_begin_line = [&](size_t align, size_t size, |
| std::string struct_name) { |
| ss << "/* " << std::setw(offset_w) << " " |
| << "align(" << std::setw(align_w) << align << ") size(" |
| << std::setw(size_w) << size << ") */ struct " << struct_name |
| << " {\n"; |
| }; |
| |
| auto print_struct_end_line = [&]() { |
| ss << "/* " |
| << std::setw(offset_w + size_w + align_w) << " " |
| << "*/ };"; |
| }; |
| |
| auto print_member_line = [&](size_t offset, size_t align, size_t size, |
| std::string s) { |
| ss << "/* offset(" << std::setw(offset_w) << offset << ") align(" |
| << std::setw(align_w) << align << ") size(" << std::setw(size_w) |
| << size << ") */ " << s << ";\n"; |
| }; |
| |
| print_struct_begin_line(st->Align(), st->Size(), |
| st->FriendlyName(builder_->Symbols())); |
| |
| for (size_t i = 0; i < st->Members().size(); ++i) { |
| auto* const m = st->Members()[i]; |
| |
| // Output field alignment padding, if any |
| auto* const prev_member = (i == 0) ? nullptr : str->Members()[i - 1]; |
| if (prev_member) { |
| uint32_t padding = |
| m->Offset() - (prev_member->Offset() + prev_member->Size()); |
| if (padding > 0) { |
| size_t padding_offset = m->Offset() - padding; |
| print_member_line(padding_offset, 1, padding, |
| "// -- implicit field alignment padding --"); |
| } |
| } |
| |
| // Output member |
| std::string member_name = member_name_of(m); |
| print_member_line(m->Offset(), m->Align(), m->Size(), |
| member_name_of(m) + " : " + type_name_of(m)); |
| } |
| |
| // Output struct size padding, if any |
| uint32_t struct_padding = st->Size() - last_member_struct_padding_offset; |
| if (struct_padding > 0) { |
| print_member_line(last_member_struct_padding_offset, 1, struct_padding, |
| "// -- implicit struct size padding --"); |
| } |
| |
| print_struct_end_line(); |
| |
| return ss.str(); |
| }; |
| |
| if (!ast::IsHostShareable(sc)) { |
| return true; |
| } |
| |
| for (size_t i = 0; i < str->Members().size(); ++i) { |
| auto* const m = str->Members()[i]; |
| uint32_t required_align = required_alignment_of(m->Type()); |
| |
| // Validate that member is at a valid byte offset |
| if (m->Offset() % required_align != 0) { |
| AddError("the offset of a struct member of type '" + type_name_of(m) + |
| "' in storage class '" + ast::ToString(sc) + |
| "' must be a multiple of " + std::to_string(required_align) + |
| " bytes, but '" + member_name_of(m) + |
| "' is currently at offset " + std::to_string(m->Offset()) + |
| ". Consider setting [[align(" + |
| std::to_string(required_align) + ")]] on this member", |
| m->Declaration()->source); |
| |
| AddNote("see layout of struct:\n" + get_struct_layout_string(str), |
| str->Declaration()->source); |
| |
| if (auto* member_str = m->Type()->As<sem::Struct>()) { |
| AddNote("and layout of struct member:\n" + |
| get_struct_layout_string(member_str), |
| member_str->Declaration()->source); |
| } |
| |
| return false; |
| } |
| |
| // For uniform buffers, validate that the number of bytes between the |
| // previous member of type struct and the current is a multiple of 16 bytes. |
| auto* const prev_member = (i == 0) ? nullptr : str->Members()[i - 1]; |
| if (prev_member && is_uniform_struct(prev_member->Type())) { |
| const uint32_t prev_to_curr_offset = m->Offset() - prev_member->Offset(); |
| if (prev_to_curr_offset % 16 != 0) { |
| AddError( |
| "uniform storage requires that the number of bytes between the " |
| "start of the previous member of type struct and the current " |
| "member be a multiple of 16 bytes, but there are currently " + |
| std::to_string(prev_to_curr_offset) + " bytes between '" + |
| member_name_of(prev_member) + "' and '" + member_name_of(m) + |
| "'. Consider setting [[align(16)]] on this member", |
| m->Declaration()->source); |
| |
| AddNote("see layout of struct:\n" + get_struct_layout_string(str), |
| str->Declaration()->source); |
| |
| auto* prev_member_str = prev_member->Type()->As<sem::Struct>(); |
| AddNote("and layout of previous member struct:\n" + |
| get_struct_layout_string(prev_member_str), |
| prev_member_str->Declaration()->source); |
| return false; |
| } |
| } |
| |
| // For uniform buffer array members, validate that array elements are |
| // aligned to 16 bytes |
| if (auto* arr = m->Type()->As<sem::Array>()) { |
| if (sc == ast::StorageClass::kUniform) { |
| // We already validated that this array member is itself aligned to 16 |
| // bytes above, so we only need to validate that stride is a multiple of |
| // 16 bytes. |
| if (arr->Stride() % 16 != 0) { |
| AddError( |
| "uniform storage requires that array elements be aligned to 16 " |
| "bytes, but array stride of '" + |
| member_name_of(m) + "' is currently " + |
| std::to_string(arr->Stride()) + |
| ". Consider setting [[stride(" + |
| std::to_string( |
| utils::RoundUp(required_align, arr->Stride())) + |
| ")]] on the array type", |
| m->Declaration()->type->source); |
| AddNote("see layout of struct:\n" + get_struct_layout_string(str), |
| str->Declaration()->source); |
| return false; |
| } |
| } |
| } |
| |
| // If member is struct, recurse |
| if (auto* str_member = m->Type()->As<sem::Struct>()) { |
| // Cache result of struct + storage class pair |
| if (valid_struct_storage_layouts_.emplace(str_member, sc).second) { |
| if (!ValidateStorageClassLayout(str_member, sc)) { |
| return false; |
| } |
| } |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateStorageClassLayout(const VariableInfo* info) { |
| if (auto* str = info->type->UnwrapRef()->As<sem::Struct>()) { |
| if (!ValidateStorageClassLayout(str, info->storage_class)) { |
| AddNote("see declaration of variable", info->declaration->source); |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateGlobalVariable(const VariableInfo* info) { |
| if (!ValidateNoDuplicateDecorations(info->declaration->decorations)) { |
| return false; |
| } |
| |
| for (auto* deco : info->declaration->decorations) { |
| if (info->declaration->is_const) { |
| if (auto* override_deco = deco->As<ast::OverrideDecoration>()) { |
| if (override_deco->has_value) { |
| uint32_t id = override_deco->value; |
| auto itr = constant_ids_.find(id); |
| if (itr != constant_ids_.end() && itr->second != info) { |
| AddError("pipeline constant IDs must be unique", deco->source); |
| AddNote("a pipeline constant with an ID of " + std::to_string(id) + |
| " was previously declared " |
| "here:", |
| ast::GetDecoration<ast::OverrideDecoration>( |
| itr->second->declaration->decorations) |
| ->source); |
| return false; |
| } |
| if (id > 65535) { |
| AddError("pipeline constant IDs must be between 0 and 65535", |
| deco->source); |
| return false; |
| } |
| } |
| } else { |
| AddError("decoration is not valid for constants", deco->source); |
| return false; |
| } |
| } else { |
| bool is_shader_io_decoration = |
| deco->IsAnyOf<ast::BuiltinDecoration, ast::InterpolateDecoration, |
| ast::InvariantDecoration, ast::LocationDecoration>(); |
| bool has_io_storage_class = |
| info->storage_class == ast::StorageClass::kInput || |
| info->storage_class == ast::StorageClass::kOutput; |
| if (!(deco->IsAnyOf<ast::BindingDecoration, ast::GroupDecoration, |
| ast::InternalDecoration>()) && |
| (!is_shader_io_decoration || !has_io_storage_class)) { |
| AddError("decoration is not valid for variables", deco->source); |
| return false; |
| } |
| } |
| } |
| |
| auto binding_point = info->declaration->BindingPoint(); |
| switch (info->storage_class) { |
| case ast::StorageClass::kUniform: |
| case ast::StorageClass::kStorage: |
| case ast::StorageClass::kUniformConstant: { |
| // https://gpuweb.github.io/gpuweb/wgsl/#resource-interface |
| // Each resource variable must be declared with both group and binding |
| // attributes. |
| if (!binding_point) { |
| AddError( |
| "resource variables require [[group]] and [[binding]] " |
| "decorations", |
| info->declaration->source); |
| return false; |
| } |
| break; |
| } |
| default: |
| if (binding_point.binding || binding_point.group) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#attribute-binding |
| // Must only be applied to a resource variable |
| AddError( |
| "non-resource variables must not have [[group]] or [[binding]] " |
| "decorations", |
| info->declaration->source); |
| return false; |
| } |
| } |
| |
| // https://gpuweb.github.io/gpuweb/wgsl/#variable-declaration |
| // The access mode always has a default, and except for variables in the |
| // storage storage class, must not be written. |
| if (info->storage_class != ast::StorageClass::kStorage && |
| info->declaration->declared_access != ast::Access::kUndefined) { |
| AddError( |
| "only variables in <storage> storage class may declare an access mode", |
| info->declaration->source); |
| return false; |
| } |
| |
| switch (info->storage_class) { |
| case ast::StorageClass::kStorage: { |
| // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables |
| // A variable in the storage storage class is a storage buffer variable. |
| // Its store type must be a host-shareable structure type with block |
| // attribute, satisfying the storage class constraints. |
| |
| auto* str = info->type->UnwrapRef()->As<sem::Struct>(); |
| |
| if (!str) { |
| AddError( |
| "variables declared in the <storage> storage class must be of a " |
| "structure type", |
| info->declaration->source); |
| return false; |
| } |
| |
| if (!str->IsBlockDecorated()) { |
| AddError( |
| "structure used as a storage buffer must be declared with the " |
| "[[block]] decoration", |
| str->Declaration()->source); |
| if (info->declaration->source.range.begin.line) { |
| AddNote("structure used as storage buffer here", |
| info->declaration->source); |
| } |
| return false; |
| } |
| break; |
| } |
| case ast::StorageClass::kUniform: { |
| // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables |
| // A variable in the uniform storage class is a uniform buffer variable. |
| // Its store type must be a host-shareable structure type with block |
| // attribute, satisfying the storage class constraints. |
| auto* str = info->type->UnwrapRef()->As<sem::Struct>(); |
| if (!str) { |
| AddError( |
| "variables declared in the <uniform> storage class must be of a " |
| "structure type", |
| info->declaration->source); |
| return false; |
| } |
| |
| if (!str->IsBlockDecorated()) { |
| AddError( |
| "structure used as a uniform buffer must be declared with the " |
| "[[block]] decoration", |
| str->Declaration()->source); |
| if (info->declaration->source.range.begin.line) { |
| AddNote("structure used as uniform buffer here", |
| info->declaration->source); |
| } |
| return false; |
| } |
| |
| for (auto* member : str->Members()) { |
| if (auto* arr = member->Type()->As<sem::Array>()) { |
| if (arr->IsRuntimeSized()) { |
| AddError( |
| "structure containing a runtime sized array " |
| "cannot be used as a uniform buffer", |
| info->declaration->source); |
| AddNote("structure is declared here", str->Declaration()->source); |
| return false; |
| } |
| } |
| } |
| |
| break; |
| } |
| default: |
| break; |
| } |
| |
| if (!info->declaration->is_const) { |
| if (!ValidateAtomicVariable(info)) { |
| return false; |
| } |
| } |
| |
| return ValidateVariable(info); |
| } |
| |
| // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types |
| // Atomic types may only be instantiated by variables in the workgroup storage |
| // class or by storage buffer variables with a read_write access mode. |
| bool Resolver::ValidateAtomicVariable(const VariableInfo* info) { |
| auto sc = info->storage_class; |
| auto access = info->access; |
| auto* type = info->type->UnwrapRef(); |
| auto source = info->declaration->type ? info->declaration->type->source |
| : info->declaration->source; |
| |
| if (type->Is<sem::Atomic>()) { |
| if (sc != ast::StorageClass::kWorkgroup) { |
| AddError( |
| "atomic variables must have <storage> or <workgroup> storage class", |
| source); |
| return false; |
| } |
| } else if (type->IsAnyOf<sem::Struct, sem::Array>()) { |
| auto found = atomic_composite_info_.find(type); |
| if (found != atomic_composite_info_.end()) { |
| if (sc != ast::StorageClass::kStorage && |
| sc != ast::StorageClass::kWorkgroup) { |
| AddError( |
| "atomic variables must have <storage> or <workgroup> storage class", |
| source); |
| AddNote("atomic sub-type of '" + |
| type->FriendlyName(builder_->Symbols()) + |
| "' is declared here", |
| found->second); |
| return false; |
| } else if (sc == ast::StorageClass::kStorage && |
| access != ast::Access::kReadWrite) { |
| AddError( |
| "atomic variables in <storage> storage class must have read_write " |
| "access mode", |
| source); |
| AddNote("atomic sub-type of '" + |
| type->FriendlyName(builder_->Symbols()) + |
| "' is declared here", |
| found->second); |
| return false; |
| } |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateVariable(const VariableInfo* info) { |
| auto* var = info->declaration; |
| auto* storage_type = info->type->UnwrapRef(); |
| |
| if (!var->is_const && !IsStorable(storage_type)) { |
| AddError(storage_type->FriendlyName(builder_->Symbols()) + |
| " cannot be used as the type of a var", |
| var->source); |
| return false; |
| } |
| |
| if (var->is_const && info->kind != VariableKind::kParameter && |
| !(storage_type->IsConstructible() || storage_type->Is<sem::Pointer>())) { |
| AddError(storage_type->FriendlyName(builder_->Symbols()) + |
| " cannot be used as the type of a let", |
| var->source); |
| return false; |
| } |
| |
| if (auto* r = storage_type->As<sem::Array>()) { |
| if (r->IsRuntimeSized()) { |
| AddError("runtime arrays may only appear as the last member of a struct", |
| var->source); |
| return false; |
| } |
| } |
| |
| if (auto* r = storage_type->As<sem::MultisampledTexture>()) { |
| if (r->dim() != ast::TextureDimension::k2d) { |
| AddError("only 2d multisampled textures are supported", var->source); |
| return false; |
| } |
| |
| if (!r->type()->UnwrapRef()->is_numeric_scalar()) { |
| AddError("texture_multisampled_2d<type>: type must be f32, i32 or u32", |
| var->source); |
| return false; |
| } |
| } |
| |
| if (storage_type->is_handle() && |
| var->declared_storage_class != ast::StorageClass::kNone) { |
| // 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 storage class decoration. The |
| // storage class will always be handle. |
| AddError("variables of type '" + info->type_name + |
| "' must not have a storage class", |
| var->source); |
| return false; |
| } |
| |
| if (IsValidationEnabled(var->decorations, |
| ast::DisabledValidation::kIgnoreStorageClass) && |
| (var->declared_storage_class == ast::StorageClass::kInput || |
| var->declared_storage_class == ast::StorageClass::kOutput)) { |
| AddError("invalid use of input/output storage class", var->source); |
| return false; |
| } |
| return true; |
| } |
| |
| bool Resolver::ValidateFunctionParameter(const ast::Function* func, |
| const VariableInfo* info) { |
| if (!ValidateVariable(info)) { |
| return false; |
| } |
| |
| for (auto* deco : info->declaration->decorations) { |
| if (!func->IsEntryPoint() && !deco->Is<ast::InternalDecoration>()) { |
| AddError( |
| "decoration is not valid for non-entry point function parameters", |
| deco->source); |
| return false; |
| } else if (!deco->IsAnyOf<ast::BuiltinDecoration, ast::InvariantDecoration, |
| ast::LocationDecoration, |
| ast::InterpolateDecoration, |
| ast::InternalDecoration>() && |
| (IsValidationEnabled( |
| info->declaration->decorations, |
| ast::DisabledValidation::kEntryPointParameter) && |
| IsValidationEnabled( |
| info->declaration->decorations, |
| ast::DisabledValidation:: |
| kIgnoreConstructibleFunctionParameter))) { |
| AddError("decoration is not valid for function parameters", deco->source); |
| return false; |
| } |
| } |
| |
| if (auto* ref = info->type->As<sem::Pointer>()) { |
| auto sc = ref->StorageClass(); |
| if (!(sc == ast::StorageClass::kFunction || |
| sc == ast::StorageClass::kPrivate || |
| sc == ast::StorageClass::kWorkgroup) && |
| IsValidationEnabled(info->declaration->decorations, |
| ast::DisabledValidation::kIgnoreStorageClass)) { |
| std::stringstream ss; |
| ss << "function parameter of pointer type cannot be in '" << sc |
| << "' storage class"; |
| AddError(ss.str(), info->declaration->source); |
| return false; |
| } |
| } |
| |
| if (IsPlain(info->type)) { |
| if (!info->type->IsConstructible() && |
| IsValidationEnabled( |
| info->declaration->decorations, |
| ast::DisabledValidation::kIgnoreConstructibleFunctionParameter)) { |
| AddError("store type of function parameter must be a constructible type", |
| info->declaration->source); |
| return false; |
| } |
| } else if (!info->type->IsAnyOf<sem::Texture, sem::Sampler, sem::Pointer>()) { |
| AddError("store type of function parameter cannot be " + |
| info->type->FriendlyName(builder_->Symbols()), |
| info->declaration->source); |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateBuiltinDecoration(const ast::BuiltinDecoration* deco, |
| const sem::Type* storage_type, |
| const bool is_input) { |
| auto* type = storage_type->UnwrapRef(); |
| const auto stage = current_function_ |
| ? current_function_->declaration->PipelineStage() |
| : ast::PipelineStage::kNone; |
| std::stringstream stage_name; |
| stage_name << stage; |
| bool is_stage_mismatch = false; |
| bool is_output = !is_input; |
| switch (deco->builtin) { |
| case ast::Builtin::kPosition: |
| if (stage != ast::PipelineStage::kNone && |
| !((is_input && stage == ast::PipelineStage::kFragment) || |
| (is_output && stage == ast::PipelineStage::kVertex))) { |
| is_stage_mismatch = true; |
| } |
| if (!(type->is_float_vector() && type->As<sem::Vector>()->Width() == 4)) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'vec4<f32>'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kGlobalInvocationId: |
| case ast::Builtin::kLocalInvocationId: |
| case ast::Builtin::kNumWorkgroups: |
| case ast::Builtin::kWorkgroupId: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kCompute && is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!(type->is_unsigned_integer_vector() && |
| type->As<sem::Vector>()->Width() == 3)) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'vec3<u32>'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kFragDepth: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kFragment && !is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::F32>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'f32'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kFrontFacing: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kFragment && is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::Bool>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'bool'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kLocalInvocationIndex: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kCompute && is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::U32>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'u32'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kVertexIndex: |
| case ast::Builtin::kInstanceIndex: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kVertex && is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::U32>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'u32'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kSampleMask: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kFragment)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::U32>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'u32'", |
| deco->source); |
| return false; |
| } |
| break; |
| case ast::Builtin::kSampleIndex: |
| if (stage != ast::PipelineStage::kNone && |
| !(stage == ast::PipelineStage::kFragment && is_input)) { |
| is_stage_mismatch = true; |
| } |
| if (!type->Is<sem::U32>()) { |
| AddError("store type of " + deco_to_str(deco) + " must be 'u32'", |
| deco->source); |
| return false; |
| } |
| break; |
| default: |
| break; |
| } |
| |
| if (is_stage_mismatch) { |
| AddError(deco_to_str(deco) + " cannot be used in " + |
| (is_input ? "input of " : "output of ") + stage_name.str() + |
| " pipeline stage", |
| deco->source); |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateInterpolateDecoration( |
| const ast::InterpolateDecoration* deco, |
| const sem::Type* storage_type) { |
| auto* type = storage_type->UnwrapRef(); |
| |
| if (type->is_integer_scalar_or_vector() && |
| deco->type != ast::InterpolationType::kFlat) { |
| AddError( |
| "interpolation type must be 'flat' for integral user-defined IO types", |
| deco->source); |
| return false; |
| } |
| |
| if (deco->type == ast::InterpolationType::kFlat && |
| deco->sampling != ast::InterpolationSampling::kNone) { |
| AddError("flat interpolation attribute must not have a sampling parameter", |
| deco->source); |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateFunction(const ast::Function* func, |
| const FunctionInfo* info) { |
| if (!ValidateNoDuplicateDefinition(func->symbol, func->source, |
| /* check_global_scope_only */ true)) { |
| return false; |
| } |
| |
| auto workgroup_deco_count = 0; |
| for (auto* deco : func->decorations) { |
| if (deco->Is<ast::WorkgroupDecoration>()) { |
| workgroup_deco_count++; |
| if (func->PipelineStage() != ast::PipelineStage::kCompute) { |
| AddError( |
| "the workgroup_size attribute is only valid for compute stages", |
| deco->source); |
| return false; |
| } |
| } else if (!deco->IsAnyOf<ast::StageDecoration, |
| ast::InternalDecoration>()) { |
| AddError("decoration is not valid for functions", deco->source); |
| return false; |
| } |
| } |
| |
| if (func->params.size() > 255) { |
| AddError("functions may declare at most 255 parameters", func->source); |
| return false; |
| } |
| |
| for (auto* param : func->params) { |
| if (!ValidateFunctionParameter(func, variable_to_info_.at(param))) { |
| return false; |
| } |
| } |
| |
| if (!info->return_type->Is<sem::Void>()) { |
| if (!info->return_type->IsConstructible()) { |
| AddError("function return type must be a constructible type", |
| func->return_type->source); |
| return false; |
| } |
| |
| if (func->body) { |
| if (!func->body->Last() || |
| !func->body->Last()->Is<ast::ReturnStatement>()) { |
| AddError("non-void function must end with a return statement", |
| func->source); |
| return false; |
| } |
| } else if (IsValidationEnabled( |
| func->decorations, |
| ast::DisabledValidation::kFunctionHasNoBody)) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "Function " << builder_->Symbols().NameFor(func->symbol) |
| << " has no body"; |
| } |
| |
| for (auto* deco : func->return_type_decorations) { |
| if (!func->IsEntryPoint()) { |
| AddError( |
| "decoration is not valid for non-entry point function return types", |
| deco->source); |
| return false; |
| } |
| if (!deco->IsAnyOf<ast::BuiltinDecoration, ast::InternalDecoration, |
| ast::LocationDecoration, ast::InterpolateDecoration, |
| ast::InvariantDecoration>() && |
| (IsValidationEnabled(info->declaration->decorations, |
| ast::DisabledValidation::kEntryPointParameter) && |
| IsValidationEnabled(info->declaration->decorations, |
| ast::DisabledValidation:: |
| kIgnoreConstructibleFunctionParameter))) { |
| AddError("decoration is not valid for entry point return types", |
| deco->source); |
| return false; |
| } |
| } |
| } |
| |
| if (func->IsEntryPoint()) { |
| if (!ValidateEntryPoint(func, info)) { |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateEntryPoint(const ast::Function* func, |
| const FunctionInfo* info) { |
| // Use a lambda to validate the entry point decorations for a type. |
| // Persistent state is used to track which builtins and locations have |
| // already been seen, in order to catch conflicts. |
| // TODO(jrprice): This state could be stored in FunctionInfo instead, and |
| // then passed to sem::Function since it would be useful there too. |
| std::unordered_set<ast::Builtin> builtins; |
| std::unordered_set<uint32_t> locations; |
| enum class ParamOrRetType { |
| kParameter, |
| kReturnType, |
| }; |
| |
| // Inner lambda that is applied to a type and all of its members. |
| auto validate_entry_point_decorations_inner = [&](const ast::DecorationList& |
| decos, |
| sem::Type* ty, |
| Source source, |
| ParamOrRetType param_or_ret, |
| bool is_struct_member) { |
| // Scan decorations for pipeline IO attributes. |
| // Check for overlap with attributes that have been seen previously. |
| const ast::Decoration* pipeline_io_attribute = nullptr; |
| const ast::InterpolateDecoration* interpolate_attribute = nullptr; |
| const ast::InvariantDecoration* invariant_attribute = nullptr; |
| for (auto* deco : decos) { |
| auto is_invalid_compute_shader_decoration = false; |
| if (auto* builtin = deco->As<ast::BuiltinDecoration>()) { |
| if (pipeline_io_attribute) { |
| AddError("multiple entry point IO attributes", deco->source); |
| AddNote("previously consumed " + deco_to_str(pipeline_io_attribute), |
| pipeline_io_attribute->source); |
| return false; |
| } |
| pipeline_io_attribute = deco; |
| |
| if (builtins.count(builtin->builtin)) { |
| AddError(deco_to_str(builtin) + |
| " attribute appears multiple times as pipeline " + |
| (param_or_ret == ParamOrRetType::kParameter ? "input" |
| : "output"), |
| func->source); |
| return false; |
| } |
| |
| if (!ValidateBuiltinDecoration( |
| builtin, ty, |
| /* is_input */ param_or_ret == ParamOrRetType::kParameter)) { |
| return false; |
| } |
| builtins.emplace(builtin->builtin); |
| } else if (auto* location = deco->As<ast::LocationDecoration>()) { |
| if (pipeline_io_attribute) { |
| AddError("multiple entry point IO attributes", deco->source); |
| AddNote("previously consumed " + deco_to_str(pipeline_io_attribute), |
| pipeline_io_attribute->source); |
| return false; |
| } |
| pipeline_io_attribute = deco; |
| |
| bool is_input = param_or_ret == ParamOrRetType::kParameter; |
| if (!ValidateLocationDecoration(location, ty, locations, source, |
| is_input)) { |
| return false; |
| } |
| } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) { |
| if (func->PipelineStage() == ast::PipelineStage::kCompute) { |
| is_invalid_compute_shader_decoration = true; |
| } else if (!ValidateInterpolateDecoration(interpolate, ty)) { |
| return false; |
| } |
| interpolate_attribute = interpolate; |
| } else if (auto* invariant = deco->As<ast::InvariantDecoration>()) { |
| if (func->PipelineStage() == ast::PipelineStage::kCompute) { |
| is_invalid_compute_shader_decoration = true; |
| } |
| invariant_attribute = invariant; |
| } |
| if (is_invalid_compute_shader_decoration) { |
| std::string input_or_output = |
| param_or_ret == ParamOrRetType::kParameter ? "inputs" : "output"; |
| AddError( |
| "decoration is not valid for compute shader " + input_or_output, |
| deco->source); |
| return false; |
| } |
| } |
| |
| if (IsValidationEnabled(decos, |
| ast::DisabledValidation::kEntryPointParameter)) { |
| if (is_struct_member && ty->Is<sem::Struct>()) { |
| AddError("nested structures cannot be used for entry point IO", source); |
| return false; |
| } |
| |
| if (!ty->Is<sem::Struct>() && !pipeline_io_attribute) { |
| std::string err = "missing entry point IO attribute"; |
| if (!is_struct_member) { |
| err += |
| (param_or_ret == ParamOrRetType::kParameter ? " on parameter" |
| : " on return type"); |
| } |
| AddError(err, source); |
| return false; |
| } |
| |
| if (pipeline_io_attribute && |
| pipeline_io_attribute->Is<ast::LocationDecoration>()) { |
| if (ty->is_integer_scalar_or_vector() && !interpolate_attribute) { |
| // TODO(crbug.com/tint/1224): Make these errors once downstream |
| // usages have caught up (no sooner than M99). |
| if (func->PipelineStage() == ast::PipelineStage::kVertex && |
| param_or_ret == ParamOrRetType::kReturnType) { |
| AddWarning( |
| "integral user-defined vertex outputs must have a flat " |
| "interpolation attribute", |
| source); |
| } |
| if (func->PipelineStage() == ast::PipelineStage::kFragment && |
| param_or_ret == ParamOrRetType::kParameter) { |
| AddWarning( |
| "integral user-defined fragment inputs must have a flat " |
| "interpolation attribute", |
| source); |
| } |
| } |
| } |
| |
| if (invariant_attribute) { |
| bool has_position = false; |
| if (pipeline_io_attribute) { |
| if (auto* builtin = |
| pipeline_io_attribute->As<ast::BuiltinDecoration>()) { |
| has_position = (builtin->builtin == ast::Builtin::kPosition); |
| } |
| } |
| if (!has_position) { |
| AddError( |
| "invariant attribute must only be applied to a position " |
| "builtin", |
| invariant_attribute->source); |
| return false; |
| } |
| } |
| } |
| return true; |
| }; |
| |
| // Outer lambda for validating the entry point decorations for a type. |
| auto validate_entry_point_decorations = [&](const ast::DecorationList& decos, |
| sem::Type* ty, Source source, |
| ParamOrRetType param_or_ret) { |
| if (!validate_entry_point_decorations_inner(decos, ty, source, param_or_ret, |
| /*is_struct_member*/ false)) { |
| return false; |
| } |
| |
| if (auto* str = ty->As<sem::Struct>()) { |
| for (auto* member : str->Members()) { |
| if (!validate_entry_point_decorations_inner( |
| member->Declaration()->decorations, member->Type(), |
| member->Declaration()->source, param_or_ret, |
| /*is_struct_member*/ true)) { |
| AddNote("while analysing entry point '" + |
| builder_->Symbols().NameFor(func->symbol) + "'", |
| func->source); |
| return false; |
| } |
| } |
| } |
| |
| return true; |
| }; |
| |
| for (auto* param : info->parameters) { |
| if (!validate_entry_point_decorations( |
| param->declaration->decorations, param->type, |
| param->declaration->source, ParamOrRetType::kParameter)) { |
| return false; |
| } |
| } |
| |
| // Clear IO sets after parameter validation. Builtin and location attributes |
| // in return types should be validated independently from those used in |
| // parameters. |
| builtins.clear(); |
| locations.clear(); |
| |
| if (!info->return_type->Is<sem::Void>()) { |
| if (!validate_entry_point_decorations(func->return_type_decorations, |
| info->return_type, func->source, |
| ParamOrRetType::kReturnType)) { |
| return false; |
| } |
| } |
| |
| if (func->PipelineStage() == ast::PipelineStage::kVertex && |
| builtins.count(ast::Builtin::kPosition) == 0) { |
| // Check module-scope variables, as the SPIR-V sanitizer generates these. |
| bool found = false; |
| for (auto* var : info->referenced_module_vars) { |
| if (auto* builtin = ast::GetDecoration<ast::BuiltinDecoration>( |
| var->declaration->decorations)) { |
| if (builtin->builtin == ast::Builtin::kPosition) { |
| found = true; |
| break; |
| } |
| } |
| } |
| if (!found) { |
| AddError( |
| "a vertex shader must include the 'position' builtin in its return " |
| "type", |
| func->source); |
| return false; |
| } |
| } |
| |
| if (func->PipelineStage() == ast::PipelineStage::kCompute) { |
| if (!ast::HasDecoration<ast::WorkgroupDecoration>(func->decorations)) { |
| AddError( |
| "a compute shader must include 'workgroup_size' in its " |
| "attributes", |
| func->source); |
| return false; |
| } |
| } |
| |
| // Validate there are no resource variable binding collisions |
| std::unordered_map<sem::BindingPoint, const ast::Variable*> binding_points; |
| for (auto* var_info : info->referenced_module_vars) { |
| if (!var_info->declaration->BindingPoint()) { |
| continue; |
| } |
| auto bp = var_info->binding_point; |
| auto res = binding_points.emplace(bp, var_info->declaration); |
| if (!res.second && |
| IsValidationEnabled(var_info->declaration->decorations, |
| ast::DisabledValidation::kBindingPointCollision) && |
| IsValidationEnabled(res.first->second->decorations, |
| ast::DisabledValidation::kBindingPointCollision)) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#resource-interface |
| // Bindings must not alias within a shader stage: two different |
| // variables in the resource interface of a given shader must not have |
| // the same group and binding values, when considered as a pair of |
| // values. |
| auto func_name = builder_->Symbols().NameFor(info->declaration->symbol); |
| AddError("entry point '" + func_name + |
| "' references multiple variables that use the " |
| "same resource binding [[group(" + |
| std::to_string(bp.group) + "), binding(" + |
| std::to_string(bp.binding) + ")]]", |
| var_info->declaration->source); |
| AddNote("first resource binding usage declared here", |
| res.first->second->source); |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::Function(const ast::Function* func) { |
| auto* info = function_infos_.Create<FunctionInfo>(func); |
| |
| if (func->IsEntryPoint()) { |
| entry_points_.emplace_back(info); |
| } |
| |
| TINT_SCOPED_ASSIGNMENT(current_function_, info); |
| |
| variable_stack_.Push(); |
| uint32_t parameter_index = 0; |
| std::unordered_map<Symbol, Source> parameter_names; |
| for (auto* param : func->params) { |
| Mark(param); |
| |
| { // Check the parameter name is unique for the function |
| auto emplaced = parameter_names.emplace(param->symbol, param->source); |
| if (!emplaced.second) { |
| auto name = builder_->Symbols().NameFor(param->symbol); |
| AddError("redefinition of parameter '" + name + "'", param->source); |
| AddNote("previous definition is here", emplaced.first->second); |
| return false; |
| } |
| } |
| |
| auto* param_info = |
| Variable(param, VariableKind::kParameter, parameter_index++); |
| if (!param_info) { |
| return false; |
| } |
| |
| for (auto* deco : param->decorations) { |
| Mark(deco); |
| } |
| if (!ValidateNoDuplicateDecorations(param->decorations)) { |
| return false; |
| } |
| |
| variable_stack_.Set(param->symbol, param_info); |
| info->parameters.emplace_back(param_info); |
| |
| if (!ApplyStorageClassUsageToType(param->declared_storage_class, |
| param_info->type, param->source)) { |
| AddNote("while instantiating parameter " + |
| builder_->Symbols().NameFor(param->symbol), |
| param->source); |
| return false; |
| } |
| if (auto* ptr = param_info->type->As<sem::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 (!ApplyStorageClassUsageToType( |
| ptr->StorageClass(), const_cast<sem::Type*>(ptr->StoreType()), |
| param->source)) { |
| AddNote("while instantiating parameter " + |
| builder_->Symbols().NameFor(param->symbol), |
| param->source); |
| return false; |
| } |
| } |
| |
| if (auto* str = param_info->type->As<sem::Struct>()) { |
| switch (func->PipelineStage()) { |
| case ast::PipelineStage::kVertex: |
| str->AddUsage(sem::PipelineStageUsage::kVertexInput); |
| break; |
| case ast::PipelineStage::kFragment: |
| str->AddUsage(sem::PipelineStageUsage::kFragmentInput); |
| break; |
| case ast::PipelineStage::kCompute: |
| str->AddUsage(sem::PipelineStageUsage::kComputeInput); |
| break; |
| case ast::PipelineStage::kNone: |
| break; |
| } |
| } |
| } |
| |
| if (auto* ty = func->return_type) { |
| info->return_type = Type(ty); |
| info->return_type_name = ty->FriendlyName(builder_->Symbols()); |
| if (!info->return_type) { |
| return false; |
| } |
| } else { |
| info->return_type = builder_->create<sem::Void>(); |
| info->return_type_name = |
| info->return_type->FriendlyName(builder_->Symbols()); |
| } |
| |
| if (auto* str = info->return_type->As<sem::Struct>()) { |
| if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, str, |
| func->source)) { |
| AddNote("while instantiating return type for " + |
| builder_->Symbols().NameFor(func->symbol), |
| func->source); |
| return false; |
| } |
| |
| switch (func->PipelineStage()) { |
| case ast::PipelineStage::kVertex: |
| str->AddUsage(sem::PipelineStageUsage::kVertexOutput); |
| break; |
| case ast::PipelineStage::kFragment: |
| str->AddUsage(sem::PipelineStageUsage::kFragmentOutput); |
| break; |
| case ast::PipelineStage::kCompute: |
| str->AddUsage(sem::PipelineStageUsage::kComputeOutput); |
| break; |
| case ast::PipelineStage::kNone: |
| break; |
| } |
| } |
| |
| if (func->body) { |
| Mark(func->body); |
| if (current_compound_statement_) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "Resolver::Function() called with a current compound statement"; |
| return false; |
| } |
| auto* sem_block = builder_->create<sem::FunctionBlockStatement>(func); |
| builder_->Sem().Add(func->body, sem_block); |
| if (!Scope(sem_block, [&] { return Statements(func->body->statements); })) { |
| return false; |
| } |
| } |
| variable_stack_.Pop(); |
| |
| for (auto* deco : func->decorations) { |
| Mark(deco); |
| } |
| if (!ValidateNoDuplicateDecorations(func->decorations)) { |
| return false; |
| } |
| |
| for (auto* deco : func->return_type_decorations) { |
| Mark(deco); |
| } |
| if (!ValidateNoDuplicateDecorations(func->return_type_decorations)) { |
| return false; |
| } |
| |
| // Set work-group size defaults. |
| for (int i = 0; i < 3; i++) { |
| info->workgroup_size[i].value = 1; |
| info->workgroup_size[i].overridable_const = nullptr; |
| } |
| |
| if (auto* workgroup = |
| ast::GetDecoration<ast::WorkgroupDecoration>(func->decorations)) { |
| auto values = workgroup->Values(); |
| auto any_i32 = false; |
| auto any_u32 = false; |
| for (int i = 0; i < 3; i++) { |
| // Each argument to this decoration can either be a literal, an |
| // identifier for a module-scope constants, or nullptr if not specified. |
| |
| auto* expr = values[i]; |
| if (!expr) { |
| // Not specified, just use the default. |
| continue; |
| } |
| |
| if (!Expression(expr)) { |
| return false; |
| } |
| |
| constexpr const char* kErrBadType = |
| "workgroup_size argument must be either literal or module-scope " |
| "constant of type i32 or u32"; |
| constexpr const char* kErrInconsistentType = |
| "workgroup_size arguments must be of the same type, either i32 " |
| "or u32"; |
| |
| auto* ty = TypeOf(expr); |
| bool is_i32 = ty->UnwrapRef()->Is<sem::I32>(); |
| bool is_u32 = ty->UnwrapRef()->Is<sem::U32>(); |
| if (!is_i32 && !is_u32) { |
| AddError(kErrBadType, expr->source); |
| return false; |
| } |
| |
| any_i32 = any_i32 || is_i32; |
| any_u32 = any_u32 || is_u32; |
| if (any_i32 && any_u32) { |
| AddError(kErrInconsistentType, expr->source); |
| return false; |
| } |
| |
| if (auto* ident = expr->As<ast::IdentifierExpression>()) { |
| // We have an identifier of a module-scope constant. |
| VariableInfo* var = variable_stack_.Get(ident->symbol); |
| if (!var || !(var->declaration->is_const)) { |
| AddError(kErrBadType, expr->source); |
| return false; |
| } |
| |
| // Capture the constant if an [[override]] attribute is present. |
| if (ast::HasDecoration<ast::OverrideDecoration>( |
| var->declaration->decorations)) { |
| info->workgroup_size[i].overridable_const = var->declaration; |
| } |
| |
| expr = var->declaration->constructor; |
| if (!expr) { |
| // No constructor means this value must be overriden by the user. |
| info->workgroup_size[i].value = 0; |
| continue; |
| } |
| } else if (!expr->Is<ast::ScalarConstructorExpression>()) { |
| AddError( |
| "workgroup_size argument must be either a literal or a " |
| "module-scope constant", |
| values[i]->source); |
| return false; |
| } |
| |
| auto val = ConstantValueOf(expr); |
| if (!val) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "could not resolve constant workgroup_size constant value"; |
| continue; |
| } |
| // Validate and set the default value for this dimension. |
| if (is_i32 ? val.Elements()[0].i32 < 1 : val.Elements()[0].u32 < 1) { |
| AddError("workgroup_size argument must be at least 1", |
| values[i]->source); |
| return false; |
| } |
| |
| info->workgroup_size[i].value = |
| is_i32 ? static_cast<uint32_t>(val.Elements()[0].i32) |
| : val.Elements()[0].u32; |
| } |
| } |
| |
| if (!ValidateFunction(func, info)) { |
| return false; |
| } |
| |
| // Register the function information _after_ processing the statements. This |
| // allows us to catch a function calling itself when determining the call |
| // information as this function doesn't exist until it's finished. |
| symbol_to_function_[func->symbol] = info; |
| function_to_info_.emplace(func, info); |
| |
| return true; |
| } |
| |
| bool Resolver::Statements(const ast::StatementList& stmts) { |
| for (auto* stmt : stmts) { |
| Mark(stmt); |
| if (!Statement(stmt)) { |
| return false; |
| } |
| } |
| if (!ValidateStatements(stmts)) { |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateStatements(const ast::StatementList& stmts) { |
| bool unreachable = false; |
| for (auto* stmt : stmts) { |
| if (unreachable) { |
| AddError("code is unreachable", stmt->source); |
| return false; |
| } |
| |
| auto* nested_stmt = stmt; |
| while (auto* block = nested_stmt->As<ast::BlockStatement>()) { |
| if (block->Empty()) { |
| break; |
| } |
| nested_stmt = block->statements.back(); |
| } |
| if (nested_stmt->IsAnyOf<ast::ReturnStatement, ast::BreakStatement, |
| ast::ContinueStatement, ast::DiscardStatement>()) { |
| unreachable = true; |
| } |
| } |
| return true; |
| } |
| |
| bool Resolver::Statement(const ast::Statement* stmt) { |
| if (stmt->Is<ast::CaseStatement>()) { |
| AddError("case statement can only be used inside a switch statement", |
| stmt->source); |
| return false; |
| } |
| if (stmt->Is<ast::ElseStatement>()) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "Resolver::Statement() encountered an Else statement. Else " |
| "statements are embedded in If statements, so should never be " |
| "encountered as top-level statements"; |
| return false; |
| } |
| |
| // Compound statements. These create their own sem::CompoundStatement |
| // bindings. |
| if (auto* b = stmt->As<ast::BlockStatement>()) { |
| return BlockStatement(b); |
| } |
| if (auto* l = stmt->As<ast::ForLoopStatement>()) { |
| return ForLoopStatement(l); |
| } |
| if (auto* l = stmt->As<ast::LoopStatement>()) { |
| return LoopStatement(l); |
| } |
| if (auto* i = stmt->As<ast::IfStatement>()) { |
| return IfStatement(i); |
| } |
| if (auto* s = stmt->As<ast::SwitchStatement>()) { |
| return SwitchStatement(s); |
| } |
| |
| // Non-Compound statements |
| sem::Statement* sem_statement = |
| builder_->create<sem::Statement>(stmt, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem_statement); |
| TINT_SCOPED_ASSIGNMENT(current_statement_, sem_statement); |
| if (auto* a = stmt->As<ast::AssignmentStatement>()) { |
| return Assignment(a); |
| } |
| if (stmt->Is<ast::BreakStatement>()) { |
| if (!sem_statement->FindFirstParent<sem::LoopBlockStatement>() && |
| !sem_statement->FindFirstParent<sem::SwitchCaseBlockStatement>()) { |
| AddError("break statement must be in a loop or switch case", |
| stmt->source); |
| return false; |
| } |
| return true; |
| } |
| if (auto* c = stmt->As<ast::CallStatement>()) { |
| if (!Expression(c->expr)) { |
| return false; |
| } |
| if (!ValidateCallStatement(c)) { |
| return false; |
| } |
| return true; |
| } |
| if (auto* c = stmt->As<ast::ContinueStatement>()) { |
| // Set if we've hit the first continue statement in our parent loop |
| if (auto* block = |
| current_block_->FindFirstParent< |
| sem::LoopBlockStatement, sem::LoopContinuingBlockStatement>()) { |
| if (auto* loop_block = block->As<sem::LoopBlockStatement>()) { |
| if (!loop_block->FirstContinue()) { |
| const_cast<sem::LoopBlockStatement*>(loop_block) |
| ->SetFirstContinue(c, loop_block->Decls().size()); |
| } |
| } else { |
| AddError("continuing blocks must not contain a continue statement", |
| stmt->source); |
| return false; |
| } |
| } else { |
| AddError("continue statement must be in a loop", stmt->source); |
| return false; |
| } |
| |
| return true; |
| } |
| if (stmt->Is<ast::DiscardStatement>()) { |
| if (auto* continuing = |
| sem_statement |
| ->FindFirstParent<sem::LoopContinuingBlockStatement>()) { |
| AddError("continuing blocks must not contain a discard statement", |
| stmt->source); |
| if (continuing != sem_statement->Parent()) { |
| AddNote("see continuing block here", continuing->Declaration()->source); |
| } |
| return false; |
| } |
| return true; |
| } |
| if (stmt->Is<ast::FallthroughStatement>()) { |
| return true; |
| } |
| if (auto* r = stmt->As<ast::ReturnStatement>()) { |
| return Return(r); |
| } |
| if (auto* v = stmt->As<ast::VariableDeclStatement>()) { |
| return VariableDeclStatement(v); |
| } |
| |
| AddError("unknown statement type for type determination: " + |
| std::string(stmt->TypeInfo().name), |
| stmt->source); |
| return false; |
| } |
| |
| bool Resolver::CaseStatement(const ast::CaseStatement* stmt) { |
| auto* sem = builder_->create<sem::SwitchCaseBlockStatement>( |
| stmt->body, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| builder_->Sem().Add(stmt->body, sem); |
| Mark(stmt->body); |
| for (auto* sel : stmt->selectors) { |
| Mark(sel); |
| } |
| return Scope(sem, [&] { return Statements(stmt->body->statements); }); |
| } |
| |
| bool Resolver::IfStatement(const ast::IfStatement* stmt) { |
| auto* sem = |
| builder_->create<sem::IfStatement>(stmt, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| return Scope(sem, [&] { |
| if (!Expression(stmt->condition)) { |
| return false; |
| } |
| |
| auto* cond_type = TypeOf(stmt->condition)->UnwrapRef(); |
| if (!cond_type->Is<sem::Bool>()) { |
| AddError("if statement condition must be bool, got " + |
| cond_type->FriendlyName(builder_->Symbols()), |
| stmt->condition->source); |
| return false; |
| } |
| |
| Mark(stmt->body); |
| auto* body = builder_->create<sem::BlockStatement>( |
| stmt->body, current_compound_statement_); |
| builder_->Sem().Add(stmt->body, body); |
| if (!Scope(body, [&] { return Statements(stmt->body->statements); })) { |
| return false; |
| } |
| |
| for (auto* else_stmt : stmt->else_statements) { |
| Mark(else_stmt); |
| if (!ElseStatement(else_stmt)) { |
| return false; |
| } |
| } |
| return true; |
| }); |
| } |
| |
| bool Resolver::ElseStatement(const ast::ElseStatement* stmt) { |
| auto* sem = |
| builder_->create<sem::ElseStatement>(stmt, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| return Scope(sem, [&] { |
| if (auto* cond = stmt->condition) { |
| if (!Expression(cond)) { |
| return false; |
| } |
| |
| auto* else_cond_type = TypeOf(cond)->UnwrapRef(); |
| if (!else_cond_type->Is<sem::Bool>()) { |
| AddError("else statement condition must be bool, got " + |
| else_cond_type->FriendlyName(builder_->Symbols()), |
| cond->source); |
| return false; |
| } |
| } |
| |
| Mark(stmt->body); |
| auto* body = builder_->create<sem::BlockStatement>( |
| stmt->body, current_compound_statement_); |
| builder_->Sem().Add(stmt->body, body); |
| return Scope(body, [&] { return Statements(stmt->body->statements); }); |
| }); |
| } |
| |
| bool Resolver::BlockStatement(const ast::BlockStatement* stmt) { |
| auto* sem = builder_->create<sem::BlockStatement>( |
| stmt->As<ast::BlockStatement>(), current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| return Scope(sem, [&] { return Statements(stmt->statements); }); |
| } |
| |
| bool Resolver::LoopStatement(const ast::LoopStatement* stmt) { |
| auto* sem = |
| builder_->create<sem::LoopStatement>(stmt, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| return Scope(sem, [&] { |
| Mark(stmt->body); |
| |
| auto* body = builder_->create<sem::LoopBlockStatement>( |
| stmt->body, current_compound_statement_); |
| builder_->Sem().Add(stmt->body, body); |
| return Scope(body, [&] { |
| if (!Statements(stmt->body->statements)) { |
| return false; |
| } |
| if (stmt->continuing) { |
| Mark(stmt->continuing); |
| if (!stmt->continuing->Empty()) { |
| auto* continuing = |
| builder_->create<sem::LoopContinuingBlockStatement>( |
| stmt->continuing, current_compound_statement_); |
| builder_->Sem().Add(stmt->continuing, continuing); |
| if (!Scope(continuing, [&] { |
| return Statements(stmt->continuing->statements); |
| })) { |
| return false; |
| } |
| } |
| } |
| return true; |
| }); |
| }); |
| } |
| |
| bool Resolver::ForLoopStatement(const ast::ForLoopStatement* stmt) { |
| auto* sem = builder_->create<sem::ForLoopStatement>( |
| stmt, current_compound_statement_); |
| builder_->Sem().Add(stmt, sem); |
| return Scope(sem, [&] { |
| if (auto* initializer = stmt->initializer) { |
| Mark(initializer); |
| if (!Statement(initializer)) { |
| return false; |
| } |
| } |
| |
| if (auto* condition = stmt->condition) { |
| if (!Expression(condition)) { |
| return false; |
| } |
| |
| if (!TypeOf(condition)->UnwrapRef()->Is<sem::Bool>()) { |
| AddError( |
| "for-loop condition must be bool, got " + TypeNameOf(condition), |
| condition->source); |
| return false; |
| } |
| } |
| |
| if (auto* continuing = stmt->continuing) { |
| Mark(continuing); |
| if (!Statement(continuing)) { |
| return false; |
| } |
| } |
| |
| Mark(stmt->body); |
| |
| auto* body = builder_->create<sem::LoopBlockStatement>( |
| stmt->body, current_compound_statement_); |
| builder_->Sem().Add(stmt->body, body); |
| return Scope(body, [&] { return Statements(stmt->body->statements); }); |
| }); |
| } |
| |
| bool Resolver::Expression(const ast::Expression* root) { |
| std::vector<const ast::Expression*> sorted; |
| if (!ast::TraverseExpressions<ast::TraverseOrder::RightToLeft>( |
| root, diagnostics_, [&](const ast::Expression* expr) { |
| Mark(expr); |
| sorted.emplace_back(expr); |
| return ast::TraverseAction::Descend; |
| })) { |
| return false; |
| } |
| |
| for (auto* expr : utils::Reverse(sorted)) { |
| bool ok = false; |
| if (auto* array = expr->As<ast::ArrayAccessorExpression>()) { |
| ok = ArrayAccessor(array); |
| } else if (auto* bin_op = expr->As<ast::BinaryExpression>()) { |
| ok = Binary(bin_op); |
| } else if (auto* bitcast = expr->As<ast::BitcastExpression>()) { |
| ok = Bitcast(bitcast); |
| } else if (auto* call = expr->As<ast::CallExpression>()) { |
| ok = Call(call); |
| } else if (auto* ctor = expr->As<ast::ConstructorExpression>()) { |
| ok = Constructor(ctor); |
| } else if (auto* ident = expr->As<ast::IdentifierExpression>()) { |
| ok = Identifier(ident); |
| } else if (auto* member = expr->As<ast::MemberAccessorExpression>()) { |
| ok = MemberAccessor(member); |
| } else if (auto* unary = expr->As<ast::UnaryOpExpression>()) { |
| ok = UnaryOp(unary); |
| } else if (expr->Is<ast::PhonyExpression>()) { |
| ok = true; // No-op |
| } else { |
| TINT_ICE(Resolver, diagnostics_) |
| << "unhandled expression type: " << expr->TypeInfo().name; |
| return false; |
| } |
| if (!ok) { |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ArrayAccessor(const ast::ArrayAccessorExpression* expr) { |
| auto* idx = expr->index; |
| auto* res = TypeOf(expr->array); |
| auto* parent_type = res->UnwrapRef(); |
| const sem::Type* ret = nullptr; |
| if (auto* arr = parent_type->As<sem::Array>()) { |
| ret = arr->ElemType(); |
| } else if (auto* vec = parent_type->As<sem::Vector>()) { |
| ret = vec->type(); |
| } else if (auto* mat = parent_type->As<sem::Matrix>()) { |
| ret = builder_->create<sem::Vector>(mat->type(), mat->rows()); |
| } else { |
| AddError("cannot index type '" + TypeNameOf(expr->array) + "'", |
| expr->source); |
| return false; |
| } |
| |
| if (!TypeOf(idx)->UnwrapRef()->IsAnyOf<sem::I32, sem::U32>()) { |
| AddError("index must be of type 'i32' or 'u32', found: '" + |
| TypeNameOf(idx) + "'", |
| idx->source); |
| return false; |
| } |
| |
| if (parent_type->Is<sem::Array>() || parent_type->Is<sem::Matrix>()) { |
| if (!res->Is<sem::Reference>()) { |
| // TODO(bclayton): expand this to allow any const_expr expression |
| // https://github.com/gpuweb/gpuweb/issues/1272 |
| auto* scalar = idx->As<ast::ScalarConstructorExpression>(); |
| if (!scalar || !scalar->literal->As<ast::IntLiteral>()) { |
| AddError("index must be signed or unsigned integer literal", |
| idx->source); |
| return false; |
| } |
| } |
| } |
| |
| // If we're extracting from a reference, we return a reference. |
| if (auto* ref = res->As<sem::Reference>()) { |
| ret = builder_->create<sem::Reference>(ret, ref->StorageClass(), |
| ref->Access()); |
| } |
| SetExprInfo(expr, ret); |
| |
| return true; |
| } |
| |
| bool Resolver::Bitcast(const ast::BitcastExpression* expr) { |
| auto* ty = Type(expr->type); |
| if (!ty) { |
| return false; |
| } |
| if (ty->Is<sem::Pointer>()) { |
| AddError("cannot cast to a pointer", expr->source); |
| return false; |
| } |
| SetExprInfo(expr, ty, expr->type->FriendlyName(builder_->Symbols())); |
| return true; |
| } |
| |
| bool Resolver::Call(const ast::CallExpression* call) { |
| Mark(call->func); |
| auto* ident = call->func; |
| auto name = builder_->Symbols().NameFor(ident->symbol); |
| |
| auto intrinsic_type = sem::ParseIntrinsicType(name); |
| if (intrinsic_type != IntrinsicType::kNone) { |
| if (!IntrinsicCall(call, intrinsic_type)) { |
| return false; |
| } |
| } else { |
| if (!FunctionCall(call)) { |
| return false; |
| } |
| } |
| |
| return ValidateCall(call); |
| } |
| |
| bool Resolver::ValidateCall(const ast::CallExpression* call) { |
| if (TypeOf(call)->Is<sem::Void>()) { |
| bool is_call_statement = false; |
| if (current_statement_) { |
| if (auto* call_stmt = |
| As<ast::CallStatement>(current_statement_->Declaration())) { |
| if (call_stmt->expr == call) { |
| is_call_statement = true; |
| } |
| } |
| } |
| if (!is_call_statement) { |
| // https://gpuweb.github.io/gpuweb/wgsl/#function-call-expr |
| // If the called function does not return a value, a function call |
| // statement should be used instead. |
| auto* ident = call->func; |
| auto name = builder_->Symbols().NameFor(ident->symbol); |
| // A function call is made to either a user declared function or an |
| // intrinsic. function_calls_ only maps CallExpression to user declared |
| // functions |
| bool is_function = function_calls_.count(call) != 0; |
| AddError((is_function ? "function" : "intrinsic") + std::string(" '") + |
| name + "' does not return a value", |
| call->source); |
| return false; |
| } |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateCallStatement(const ast::CallStatement*) { |
| return true; |
| } |
| |
| bool Resolver::IntrinsicCall(const ast::CallExpression* call, |
| sem::IntrinsicType intrinsic_type) { |
| std::vector<const sem::Type*> arg_tys; |
| arg_tys.reserve(call->args.size()); |
| for (auto* expr : call->args) { |
| arg_tys.emplace_back(TypeOf(expr)); |
| } |
| |
| auto* result = |
| intrinsic_table_->Lookup(intrinsic_type, arg_tys, call->source); |
| if (!result) { |
| return false; |
| } |
| |
| if (result->IsDeprecated()) { |
| AddWarning("use of deprecated intrinsic", call->source); |
| } |
| |
| auto* out = builder_->create<sem::Call>(call, result, current_statement_); |
| builder_->Sem().Add(call, out); |
| SetExprInfo(call, result->ReturnType()); |
| |
| current_function_->intrinsic_calls.emplace_back( |
| IntrinsicCallInfo{call, result}); |
| |
| if (IsTextureIntrinsic(intrinsic_type) && |
| !ValidateTextureIntrinsicFunction(call, out)) { |
| return false; |
| } |
| |
| return true; |
| } |
| |
| bool Resolver::ValidateTextureIntrinsicFunction( |
| const ast::CallExpression* ast_call, |
| const sem::Call* sem_call) { |
| auto* intrinsic = sem_call->Target()->As<sem::Intrinsic>(); |
| if (!intrinsic) { |
| return false; |
| } |
| std::string func_name = intrinsic->str(); |
| auto& signature = intrinsic->Signature(); |
| auto index = signature.IndexOf(sem::ParameterUsage::kOffset); |
| if (index > -1) { |
| auto* param = ast_call->args[index]; |
| if (param->Is<ast::TypeConstructorExpression>()) { |
| auto values = ConstantValueOf(param); |
| if (!values.IsValid()) { |
| AddError( |
| "'" + func_name + "' offset parameter must be a const_expression", |
| param->source); |
| return false; |
| } |
| if (!values.Type()->Is<sem::Vector>() || |
| !values.ElementType()->Is<sem::I32>()) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "failed to resolve '" + func_name + "' offset parameter type"; |
| return false; |
| } |
| for (auto offset : values.Elements()) { |
| auto offset_value = offset.i32; |
| if (offset_value < -8 || offset_value > 7) { |
| AddError("each offset component of '" + func_name + |
| "' must be at least -8 and at most 7. " |
| "found: '" + |
| std::to_string(offset_value) + "'", |
| param->source); |
| return false; |
| } |
| } |
| } else { |
| AddError( |
| "'" + func_name + "' offset parameter must be a const_expression", |
| param->source); |
| return false; |
| } |
| } |
| return true; |
| } |
| |
| bool Resolver::FunctionCall(const ast::CallExpression* call) { |
| auto* ident = call->func; |
| auto name = builder_->Symbols().NameFor(ident->symbol); |
| |
| auto callee_func_it = symbol_to_function_.find(ident->symbol); |
| if (callee_func_it == symbol_to_function_.end()) { |
| if (current_function_ && |
| current_function_->declaration->symbol == ident->symbol) { |
| AddError("recursion is not permitted. '" + name + |
| "' attempted to call itself.", |
| call->source); |
| } else { |
| AddError("unable to find called function: " + name, call->source); |
| } |
| return false; |
| } |
| auto* callee_func = callee_func_it->second; |
| |
| if (current_function_) { |
| callee_func->callsites.push_back(call); |
| |
| // Note: Requires called functions to be resolved first. |
| // This is currently guaranteed as functions must be declared before |
| // use. |
| current_function_->transitive_calls.add(callee_func); |
| for (auto* transitive_call : callee_func->transitive_calls) { |
| current_function_->transitive_calls.add(transitive_call); |
| } |
| |
| // We inherit any referenced variables from the callee. |
| for (auto* var : callee_func->referenced_module_vars) { |
| set_referenced_from_function_if_needed(var, false); |
| } |
| } |
| |
| function_calls_.emplace(call, |
| FunctionCallInfo{callee_func, current_statement_}); |
| SetExprInfo(call, callee_func->return_type, callee_func->return_type_name); |
| |
| if (!ValidateFunctionCall(call, callee_func)) { |
| return false; |
| } |
| return true; |
| } |
| |
| bool Resolver::ValidateFunctionCall(const ast::CallExpression* call, |
| const FunctionInfo* target) { |
| auto* ident = call->func; |
| auto name = builder_->Symbols().NameFor(ident->symbol); |
| |
| if (target->declaration->IsEntryPoint()) { |
| // https://www.w3.org/TR/WGSL/#function-restriction |
| // An entry point must never be the target of a function call. |
| AddError("entry point functions cannot be the target of a function call", |
| call->source); |
| return false; |
| } |
| |
| if (call->args.size() != target->parameters.size()) { |
| bool more = call->args.size() > target->parameters.size(); |
| AddError("too " + (more ? std::string("many") : std::string("few")) + |
| " arguments in call to '" + name + "', expected " + |
| std::to_string(target->parameters.size()) + ", got " + |
| std::to_string(call->args.size()), |
| call->source); |
| return false; |
| } |
| |
| for (size_t i = 0; i < call->args.size(); ++i) { |
| const VariableInfo* param = target->parameters[i]; |
| const ast::Expression* arg_expr = call->args[i]; |
| auto* arg_type = TypeOf(arg_expr)->UnwrapRef(); |
| |
| if (param->type != arg_type) { |
| AddError("type mismatch for argument " + std::to_string(i + 1) + |
| " in call to '" + name + "', expected '" + |
| param->type->FriendlyName(builder_->Symbols()) + "', got '" + |
| arg_type->FriendlyName(builder_->Symbols()) + "'", |
| arg_expr->source); |
| return false; |
| } |
| |
| if (param->declaration->type->Is<ast::Pointer>()) { |
| auto is_valid = false; |
| if (auto* ident_expr = arg_expr->As<ast::IdentifierExpression>()) { |
| VariableInfo* var = variable_stack_.Get(ident_expr->symbol); |
| if (!var) { |
| TINT_ICE(Resolver, diagnostics_) << "failed to resolve identifier"; |
| return false; |
| } |
| if (var->kind == VariableKind::kParameter) { |
| is_valid = true; |
| } |
| } else if (auto* unary = arg_expr->As<ast::UnaryOpExpression>()) { |
| if (unary->op == ast::UnaryOp::kAddressOf) { |
| if (auto* ident_unary = |
| unary->expr->As<ast::IdentifierExpression>()) { |
| VariableInfo* var = variable_stack_.Get(ident_unary->symbol); |
| if (!var) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "failed to resolve identifier"; |
| return false; |
| } |
| if (var->declaration->is_const) { |
| TINT_ICE(Resolver, diagnostics_) |
| << "Resolver::FunctionCall() encountered an address-of " |
| "expression of a constant identifier expression"; |
| return false; |
| } |
| is_valid = true; |
| } |
| } |
| } |
| |
| if (!is_valid && |
| IsValidationEnabled( |
| param->declaration->decorations, |
| ast::DisabledValidation::kIgnoreInvalidPointerArgument)) { |
| AddError( |
| "expected an address-of expression of a variable identifier " |
| "expression or a function parameter", |
| arg_expr->source); |
| return false; |
| } |
| } |
| } |
| return true; |
| } |
| |
| bool Resolver::Constructor(const ast::ConstructorExpression* expr) { |
| if (auto* type_ctor = expr->As<ast::TypeConstructorExpression>()) { |
| auto* type = Type(type_ctor->type); |
| if (!type) { |
| return false; |
| } |
| |
| auto type_name = type_ctor->type->FriendlyName(builder_->Symbols()); |
| |
| // Now that the argument types have been determined, make sure that they |
| // obey the constructor type rules laid out in |
| // https://gpuweb.github.io/gpuweb/wgsl.html#type-constructor-expr. |
| bool ok = true; |
| if (auto* vec_type = type->As<sem::Vector>()) { |
| ok = ValidateVectorConstructor(type_ctor, vec_type, type_name); |
| } else if (auto* mat_type = type->As<sem::Matrix>()) { |
| ok = ValidateMatrixConstructor(type_ctor, mat_type, type_name); |
| } else if (type->is_scalar()) { |
| ok = ValidateScalarConstructor(type_ctor, type, type_name); |
| } else if (auto* arr_type = type->As<sem::Array>()) { |
| ok = ValidateArrayConstructor(type_ctor, arr_type); |
| } else if (auto* struct_type = type->As<sem::Struct>()) { |
| ok = ValidateStructureConstructor(type_ctor, struct_type); |
| } else { |
| AddError("type is not constructible", type_ctor->source); |
| return false; |
|