diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index dd679c7f02..fbed83aca3 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1646,7 +1646,6 @@ if (tint_build_unittests) { "reader/wgsl/classify_template_args_test.cc", "reader/wgsl/lexer_test.cc", "reader/wgsl/parser_impl_additive_expression_test.cc", - "reader/wgsl/parser_impl_address_space_test.cc", "reader/wgsl/parser_impl_argument_expression_list_test.cc", "reader/wgsl/parser_impl_assignment_stmt_test.cc", "reader/wgsl/parser_impl_bitwise_expression_test.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 67ca5e9087..4b87d7b4a8 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -1092,7 +1092,6 @@ if(TINT_BUILD_TESTS) reader/wgsl/parser_impl_singular_expression_test.cc reader/wgsl/parser_impl_statement_test.cc reader/wgsl/parser_impl_statements_test.cc - reader/wgsl/parser_impl_address_space_test.cc reader/wgsl/parser_impl_struct_body_decl_test.cc reader/wgsl/parser_impl_struct_decl_test.cc reader/wgsl/parser_impl_struct_attribute_decl_test.cc diff --git a/src/tint/ast/var.cc b/src/tint/ast/var.cc index b748381470..bccdbccfad 100644 --- a/src/tint/ast/var.cc +++ b/src/tint/ast/var.cc @@ -25,8 +25,8 @@ Var::Var(ProgramID pid, const Source& src, const Identifier* n, Type ty, - type::AddressSpace address_space, - type::Access access, + const Expression* address_space, + const Expression* access, const Expression* init, utils::VectorRef attrs) : Base(pid, nid, src, n, ty, init, std::move(attrs)), @@ -45,10 +45,11 @@ const Var* Var::Clone(CloneContext* ctx) const { auto src = ctx->Clone(source); auto* n = ctx->Clone(name); auto ty = ctx->Clone(type); + auto* address_space = ctx->Clone(declared_address_space); + auto* access = ctx->Clone(declared_access); auto* init = ctx->Clone(initializer); auto attrs = ctx->Clone(attributes); - return ctx->dst->create(src, n, ty, declared_address_space, declared_access, init, - std::move(attrs)); + return ctx->dst->create(src, n, ty, address_space, access, init, std::move(attrs)); } } // namespace tint::ast diff --git a/src/tint/ast/var.h b/src/tint/ast/var.h index e3537d1e32..006d8de385 100644 --- a/src/tint/ast/var.h +++ b/src/tint/ast/var.h @@ -56,8 +56,8 @@ class Var final : public Castable { const Source& source, const Identifier* name, Type type, - type::AddressSpace declared_address_space, - type::Access declared_access, + const Expression* declared_address_space, + const Expression* declared_access, const Expression* initializer, utils::VectorRef attributes); @@ -77,10 +77,10 @@ class Var final : public Castable { const Var* Clone(CloneContext* ctx) const override; /// The declared address space - const type::AddressSpace declared_address_space; + const Expression* const declared_address_space = nullptr; /// The declared access control - const type::Access declared_access; + const Expression* const declared_access = nullptr; }; /// A list of `var` declarations diff --git a/src/tint/ast/variable_test.cc b/src/tint/ast/variable_test.cc index 88bcdc46a5..dcbf1e5a24 100644 --- a/src/tint/ast/variable_test.cc +++ b/src/tint/ast/variable_test.cc @@ -28,7 +28,8 @@ TEST_F(VariableTest, Creation) { auto* v = Var("my_var", ty.i32(), type::AddressSpace::kFunction); CheckIdentifier(Symbols(), v->name, "my_var"); - EXPECT_EQ(v->declared_address_space, type::AddressSpace::kFunction); + CheckIdentifier(Symbols(), v->declared_address_space, "function"); + EXPECT_EQ(v->declared_access, nullptr); CheckIdentifier(Symbols(), v->type, "i32"); EXPECT_EQ(v->source.range.begin.line, 0u); EXPECT_EQ(v->source.range.begin.column, 0u); @@ -41,7 +42,7 @@ TEST_F(VariableTest, CreationWithSource) { ty.f32(), type::AddressSpace::kPrivate, utils::Empty); CheckIdentifier(Symbols(), v->name, "i"); - EXPECT_EQ(v->declared_address_space, type::AddressSpace::kPrivate); + CheckIdentifier(Symbols(), v->declared_address_space, "private"); CheckIdentifier(Symbols(), v->type, "f32"); EXPECT_EQ(v->source.range.begin.line, 27u); EXPECT_EQ(v->source.range.begin.column, 4u); @@ -54,7 +55,7 @@ TEST_F(VariableTest, CreationEmpty) { ty.i32(), type::AddressSpace::kWorkgroup, utils::Empty); CheckIdentifier(Symbols(), v->name, "a_var"); - EXPECT_EQ(v->declared_address_space, type::AddressSpace::kWorkgroup); + CheckIdentifier(Symbols(), v->declared_address_space, "workgroup"); CheckIdentifier(Symbols(), v->type, "i32"); EXPECT_EQ(v->source.range.begin.line, 27u); EXPECT_EQ(v->source.range.begin.column, 4u); diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index ee280106d9..f55b12cd6e 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -80,9 +80,9 @@ enum address_space { uniform storage push_constant + __in + __out @internal handle - @internal in - @internal out } // https://gpuweb.github.io/gpuweb/wgsl/#memory-access-mode diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h index 1c04faf85c..f89a30baee 100644 --- a/src/tint/program_builder.h +++ b/src/tint/program_builder.h @@ -219,24 +219,34 @@ class ProgramBuilder { /// constructing an ast::Var. struct VarOptions { template - explicit VarOptions(ARGS&&... args) { - (Set(std::forward(args)), ...); + explicit VarOptions(ProgramBuilder& b, ARGS&&... args) { + (Set(b, std::forward(args)), ...); } ~VarOptions(); ast::Type type; - type::AddressSpace address_space = type::AddressSpace::kUndefined; - type::Access access = type::Access::kUndefined; + const ast::Expression* address_space = nullptr; + const ast::Expression* access = nullptr; const ast::Expression* initializer = nullptr; utils::Vector attributes; private: - void Set(ast::Type t) { type = t; } - void Set(type::AddressSpace addr_space) { address_space = addr_space; } - void Set(type::Access ac) { access = ac; } - void Set(const ast::Expression* c) { initializer = c; } - void Set(utils::VectorRef l) { attributes = std::move(l); } - void Set(const ast::Attribute* a) { attributes.Push(a); } + void Set(ProgramBuilder&, ast::Type t) { type = t; } + void Set(ProgramBuilder& b, type::AddressSpace addr_space) { + if (addr_space != type::AddressSpace::kUndefined) { + address_space = b.Expr(addr_space); + } + } + void Set(ProgramBuilder& b, type::Access ac) { + if (ac != type::Access::kUndefined) { + access = b.Expr(ac); + } + } + void Set(ProgramBuilder&, const ast::Expression* c) { initializer = c; } + void Set(ProgramBuilder&, utils::VectorRef l) { + attributes = std::move(l); + } + void Set(ProgramBuilder&, const ast::Attribute* a) { attributes.Push(a); } }; /// LetOptions is a helper for accepting an arbitrary number of order independent options for @@ -2055,9 +2065,9 @@ class ProgramBuilder { /// @param name the variable name /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: - /// * ast::Type* - specifies the variable type - /// * type::AddressSpace - specifies the variable address space - /// * type::Access - specifies the variable's access control + /// * ast::Type - specifies the variable's type + /// * type::AddressSpace - specifies the variable's address space + /// * type::Access - specifies the variable's access control /// * ast::Expression* - specifies the variable's initializer expression /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. @@ -2072,16 +2082,16 @@ class ProgramBuilder { /// @param name the variable name /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: - /// * ast::Type* - specifies the variable type - /// * type::AddressSpace - specifies the variable address space - /// * type::Access - specifies the variable's access control + /// * ast::Type - specifies the variable's type + /// * type::AddressSpace - specifies the variable's address space + /// * type::Access - specifies the variable's access control /// * ast::Expression* - specifies the variable's initializer expression /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns a `ast::Var` with the given name, address_space and type template const ast::Var* Var(const Source& source, NAME&& name, OPTIONS&&... options) { - VarOptions opts(std::forward(options)...); + VarOptions opts(*this, std::forward(options)...); return create(source, Ident(std::forward(name)), opts.type, opts.address_space, opts.access, opts.initializer, std::move(opts.attributes)); @@ -2091,8 +2101,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Identifier* - specifies the variable type - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Const` with the given name, type and additional options @@ -2106,8 +2115,8 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Identifier* - specifies the variable type - /// * ast::Type* - specifies the variable type + /// * ast::Identifier* - specifies the variable's type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Const` with the given name, type and additional options @@ -2122,7 +2131,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Let` with the given name, type and additional options @@ -2136,7 +2145,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Let` with the given name, type and additional options @@ -2175,7 +2184,7 @@ class ProgramBuilder { /// @param name the variable name /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * type::AddressSpace - specifies the variable address space /// * type::Access - specifies the variable's access control /// * ast::Expression* - specifies the variable's initializer expression @@ -2192,7 +2201,7 @@ class ProgramBuilder { /// @param name the variable name /// @param options the extra options passed to the ast::Var initializer /// Can be any of the following, in any order: - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * type::AddressSpace - specifies the variable address space /// * type::Access - specifies the variable's access control /// * ast::Expression* - specifies the variable's initializer expression @@ -2211,7 +2220,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Const initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Const` with the given name, type and additional options, which is @@ -2226,7 +2235,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Const initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Const` with the given name, type and additional options, which is @@ -2242,7 +2251,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Override initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Override` with the given name, type and additional options, which is @@ -2257,7 +2266,7 @@ class ProgramBuilder { /// @param options the extra options passed to the ast::Override initializer /// Can be any of the following, in any order: /// * ast::Expression* - specifies the variable's initializer expression (required) - /// * ast::Type* - specifies the variable type + /// * ast::Type - specifies the variable's type /// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector) /// Note that non-repeatable arguments of the same type will use the last argument's value. /// @returns an `ast::Override` with the given name, type and additional options, which is diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc index 046780fec1..513b7f3a81 100644 --- a/src/tint/reader/spirv/function.cc +++ b/src/tint/reader/spirv/function.cc @@ -2522,7 +2522,8 @@ bool FunctionEmitter::EmitFunctionVariables() { } } auto* var = parser_impl_.MakeVar(inst.result_id(), type::AddressSpace::kUndefined, - var_store_type, initializer, AttributeList{}); + type::Access::kUndefined, var_store_type, initializer, + AttributeList{}); auto* var_decl_stmt = create(Source{}, var); AddStatement(var_decl_stmt); auto* var_type = ty_.Reference(var_store_type, type::AddressSpace::kUndefined); @@ -3367,8 +3368,9 @@ bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info, // no need to remap pointer properties. auto* store_type = parser_impl_.ConvertType(def_inst->type_id()); AddStatement(create( - Source{}, parser_impl_.MakeVar(id, type::AddressSpace::kUndefined, store_type, nullptr, - AttributeList{}))); + Source{}, + parser_impl_.MakeVar(id, type::AddressSpace::kUndefined, type::Access::kUndefined, + store_type, nullptr, AttributeList{}))); auto* type = ty_.Reference(store_type, type::AddressSpace::kUndefined); identifier_types_.emplace(id, type); } @@ -4835,9 +4837,8 @@ DefInfo::Pointer FunctionEmitter::GetPointerInfo(uint32_t id) { // either variables or function parameters. switch (opcode(inst)) { case spv::Op::OpVariable: { - if (const auto* module_var = parser_impl_.GetModuleVariable(id)) { - return DefInfo::Pointer{module_var->declared_address_space, - module_var->declared_access}; + if (auto v = parser_impl_.GetModuleVariable(id); v.var) { + return DefInfo::Pointer{v.address_space, v.access}; } // Local variables are always Function storage class, with default // access mode. diff --git a/src/tint/reader/spirv/parser_impl.cc b/src/tint/reader/spirv/parser_impl.cc index 5707bf14e6..443b28f357 100644 --- a/src/tint/reader/spirv/parser_impl.cc +++ b/src/tint/reader/spirv/parser_impl.cc @@ -1504,12 +1504,15 @@ bool ParserImpl::EmitModuleScopeVariables() { // here.) ast_initializer = MakeConstantExpression(var.GetSingleWordInOperand(1)).expr; } - auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_store_type, ast_initializer, - utils::Empty); + auto ast_access = VarAccess(ast_store_type, ast_address_space); + auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_access, ast_store_type, + ast_initializer, utils::Empty); // TODO(dneto): initializers (a.k.a. initializer expression) if (ast_var) { builder_.AST().AddGlobalVariable(ast_var); - module_variable_.GetOrCreate(var.result_id(), [ast_var] { return ast_var; }); + module_variable_.GetOrCreate(var.result_id(), [&] { + return ModuleVariable{ast_var, ast_address_space, ast_access}; + }); } } @@ -1536,14 +1539,16 @@ bool ParserImpl::EmitModuleScopeVariables() { << init->PrettyPrint(); } } - auto* ast_var = - MakeVar(builtin_position_.per_vertex_var_id, - enum_converter_.ToAddressSpace(builtin_position_.storage_class), - ConvertType(builtin_position_.position_member_type_id), ast_initializer, {}); + auto storage_type = ConvertType(builtin_position_.position_member_type_id); + auto ast_address_space = enum_converter_.ToAddressSpace(builtin_position_.storage_class); + auto ast_access = VarAccess(storage_type, ast_address_space); + auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space, ast_access, + storage_type, ast_initializer, {}); builder_.AST().AddGlobalVariable(ast_var); - module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id, - [ast_var] { return ast_var; }); + module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id, [&] { + return ModuleVariable{ast_var, ast_address_space}; + }); } return success_; } @@ -1571,8 +1576,23 @@ const spvtools::opt::analysis::IntConstant* ParserImpl::GetArraySize(uint32_t va return size->AsIntConstant(); } +type::Access ParserImpl::VarAccess(const Type* storage_type, type::AddressSpace address_space) { + if (address_space != type::AddressSpace::kStorage) { + return type::Access::kUndefined; + } + + bool read_only = false; + if (auto* tn = storage_type->As()) { + read_only = read_only_struct_types_.count(tn->name) > 0; + } + + // Apply the access(read) or access(read_write) modifier. + return read_only ? type::Access::kRead : type::Access::kReadWrite; +} + const ast::Var* ParserImpl::MakeVar(uint32_t id, type::AddressSpace address_space, + type::Access access, const Type* storage_type, const ast::Expression* initializer, AttributeList decorations) { @@ -1581,17 +1601,6 @@ const ast::Var* ParserImpl::MakeVar(uint32_t id, return nullptr; } - type::Access access = type::Access::kUndefined; - if (address_space == type::AddressSpace::kStorage) { - bool read_only = false; - if (auto* tn = storage_type->As()) { - read_only = read_only_struct_types_.count(tn->name) > 0; - } - - // Apply the access(read) or access(read_write) modifier. - access = read_only ? type::Access::kRead : type::Access::kReadWrite; - } - // Handle variables (textures and samplers) are always in the handle // address space, so we don't mention the address space. if (address_space == type::AddressSpace::kHandle) { diff --git a/src/tint/reader/spirv/parser_impl.h b/src/tint/reader/spirv/parser_impl.h index 5d6daa4380..0c77a02849 100644 --- a/src/tint/reader/spirv/parser_impl.h +++ b/src/tint/reader/spirv/parser_impl.h @@ -420,10 +420,17 @@ class ParserImpl : Reader { /// @returns a list of SPIR-V decorations. DecorationList GetMemberPipelineDecorations(const Struct& struct_type, int member_index); + /// @param storage_type the 'var' storage type + /// @param address_space the 'var' address space + /// @returns the access mode for a 'var' declaration with the given storage type and address + /// space. + type::Access VarAccess(const Type* storage_type, type::AddressSpace address_space); + /// Creates an AST 'var' node for a SPIR-V ID, including any attached decorations, unless it's /// an ignorable builtin variable. /// @param id the SPIR-V result ID /// @param address_space the address space, which cannot be type::AddressSpace::kUndefined + /// @param access the access /// @param storage_type the storage type of the variable /// @param initializer the variable initializer /// @param decorations the variable decorations @@ -431,6 +438,7 @@ class ParserImpl : Reader { /// in the error case const ast::Var* MakeVar(uint32_t id, type::AddressSpace address_space, + type::Access access, const Type* storage_type, const ast::Expression* initializer, AttributeList decorations); @@ -659,13 +667,23 @@ class ParserImpl : Reader { /// error const Type* GetHandleTypeForSpirvHandle(const spvtools::opt::Instruction& obj); + /// ModuleVariable describes a module scope variable + struct ModuleVariable { + /// The AST variable node. + const ast::Var* var = nullptr; + /// The address space of the var + type::AddressSpace address_space = type::AddressSpace::kUndefined; + /// The access mode of the var + type::Access access = type::Access::kUndefined; + }; + /// Returns the AST variable for the SPIR-V ID of a module-scope variable, /// or null if there isn't one. /// @param id a SPIR-V ID /// @returns the AST variable or null. - const ast::Var* GetModuleVariable(uint32_t id) { + ModuleVariable GetModuleVariable(uint32_t id) { auto entry = module_variable_.Find(id); - return entry ? *entry : nullptr; + return entry ? *entry : ModuleVariable{}; } /// Returns the channel component type corresponding to the given image @@ -885,7 +903,7 @@ class ParserImpl : Reader { std::unordered_map handle_type_; /// Maps the SPIR-V ID of a module-scope variable to its AST variable. - utils::Hashmap module_variable_; + utils::Hashmap module_variable_; // Set of symbols of declared type that have been added, used to avoid // adding duplicates. diff --git a/src/tint/reader/wgsl/classify_template_args.cc b/src/tint/reader/wgsl/classify_template_args.cc index 0ccef952e3..79d39d2228 100644 --- a/src/tint/reader/wgsl/classify_template_args.cc +++ b/src/tint/reader/wgsl/classify_template_args.cc @@ -70,8 +70,8 @@ void ClassifyTemplateArguments(std::vector& tokens) { for (size_t i = 0; i < count - 1; i++) { switch (tokens[i].type()) { - // + all type / builtin keywords that will become identifiers. case Token::Type::kIdentifier: + case Token::Type::kVar: case Token::Type::kBitcast: { auto& next = tokens[i + 1]; if (next.type() == Token::Type::kLessThan) { diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc index c50335343e..c9a8576e4c 100644 --- a/src/tint/reader/wgsl/parser_impl.cc +++ b/src/tint/reader/wgsl/parser_impl.cc @@ -204,23 +204,6 @@ ParserImpl::FunctionHeader::~FunctionHeader() = default; ParserImpl::FunctionHeader& ParserImpl::FunctionHeader::operator=(const FunctionHeader& rhs) = default; -ParserImpl::VarDeclInfo::VarDeclInfo() = default; - -ParserImpl::VarDeclInfo::VarDeclInfo(const VarDeclInfo&) = default; - -ParserImpl::VarDeclInfo::VarDeclInfo(Source source_in, - std::string name_in, - type::AddressSpace address_space_in, - type::Access access_in, - ast::Type type_in) - : source(std::move(source_in)), - name(std::move(name_in)), - address_space(address_space_in), - access(access_in), - type(type_in) {} - -ParserImpl::VarDeclInfo::~VarDeclInfo() = default; - ParserImpl::ParserImpl(Source::File const* file) : file_(file) {} ParserImpl::~ParserImpl() = default; @@ -614,13 +597,13 @@ Maybe ParserImpl::global_variable_decl(AttributeList& attr TINT_DEFER(attrs.Clear()); - return builder_.Var(decl->source, // source - decl->name, // symbol - decl->type, // type - decl->address_space, // address space - decl->access, // access control - initializer, // initializer - std::move(attrs)); // attributes + return builder_.create(decl->source, // source + builder_.Ident(decl->name), // symbol + decl->type, // type + decl->address_space, // address space + decl->access, // access control + initializer, // initializer + std::move(attrs)); // attributes } // global_constant_decl : @@ -756,37 +739,28 @@ Expect ParserImpl::expect_ident_with_type_specifier return expect_ident_with_optional_type_specifier(use, /* allow_inferred */ false); } -// access_mode -// : 'read' -// | 'write' -// | 'read_write' -Expect ParserImpl::expect_access_mode(std::string_view use) { - return expect_enum("access control", type::ParseAccess, type::kAccessStrings, use); -} - // variable_qualifier -// : LESS_THAN address_spaces (COMMA access_mode)? GREATER_THAN +// : _template_args_start expression (COMMA expression)? _template_args_end Maybe ParserImpl::variable_qualifier() { - if (!peek_is(Token::Type::kLessThan)) { + if (!peek_is(Token::Type::kTemplateArgsLeft) && !peek_is(Token::Type::kLessThan)) { + // Note: kLessThan will give a sensible error at expect_template_arg_block() return Failure::kNoMatch; } auto* use = "variable declaration"; - auto vq = expect_lt_gt_block(use, [&]() -> Expect { - auto source = make_source_range(); - auto sc = expect_address_space(use); - if (sc.errored) { + auto vq = expect_template_arg_block(use, [&]() -> Expect { + auto address_space = expect_expression("'var' address space"); + if (address_space.errored) { return Failure::kErrored; } if (match(Token::Type::kComma)) { - auto ac = expect_access_mode(use); - if (ac.errored) { + auto access = expect_expression("'var' access mode"); + if (access.errored) { return Failure::kErrored; } - return VariableQualifier{sc.value, ac.value}; + return VariableQualifier{address_space.value, access.value}; } - return Expect{VariableQualifier{sc.value, type::Access::kUndefined}, - source}; + return VariableQualifier{address_space.value}; }); if (vq.errored) { @@ -900,18 +874,6 @@ Expect ParserImpl::expect_type(std::string_view use) { return type.value; } -// address_space -// : 'function' -// | 'private' -// | 'workgroup' -// | 'uniform' -// | 'storage' -// -// Note, we also parse `push_constant` from the experimental extension -Expect ParserImpl::expect_address_space(std::string_view use) { - return expect_enum("address space", type::ParseAddressSpace, type::kAddressSpaceStrings, use); -} - // struct_decl // : STRUCT IDENT struct_body_decl Maybe ParserImpl::struct_decl() { @@ -1519,12 +1481,13 @@ Maybe ParserImpl::variable_statement() { initializer = initializer_expr.value; } - auto* var = builder_.Var(decl_source, // source - decl->name, // symbol - decl->type, // type - decl->address_space, // address space - decl->access, // access control - initializer); // initializer + auto* var = builder_.create(decl_source, // source + builder_.Ident(decl->name), // symbol + decl->type, // type + decl->address_space, // address space + decl->access, // access control + initializer, // initializer + utils::Empty); // attributes return create(var->source, var); } @@ -2520,7 +2483,7 @@ Expect ParserImpl::expect_relational_expression_post_una return create(tok_op.source(), op, lhs, rhs.value); } -Expect ParserImpl::expect_expression() { +Expect ParserImpl::expect_expression(std::string_view use) { auto& t = peek(); auto expr = expression(); if (expr.errored) { @@ -2529,7 +2492,7 @@ Expect ParserImpl::expect_expression() { if (expr.matched) { return expr.value; } - return add_error(t, "expected expression"); + return add_error(t, "expected expression for " + std::string(use)); } Expect> ParserImpl::expect_expression_list( @@ -2537,7 +2500,7 @@ Expect> ParserImpl::expect_expression_l Token::Type terminator) { utils::Vector exprs; while (continue_parsing()) { - auto expr = expect_expression(); + auto expr = expect_expression(use); if (expr.errored) { return Failure::kErrored; } diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h index f692d777b2..64618b3704 100644 --- a/src/tint/reader/wgsl/parser_impl.h +++ b/src/tint/reader/wgsl/parser_impl.h @@ -148,10 +148,9 @@ class ParserImpl { /// Constructor for a successful parse. /// @param val the result value of the parse - /// @param s the optional source of the value template - inline Maybe(U&& val, const Source& s = {}) // NOLINT - : value(std::forward(val)), source(s), matched(true) {} + inline Maybe(U&& val) // NOLINT + : value(std::forward(val)), matched(true) {} /// Constructor for parse error state. inline Maybe(Failure::Errored) : errored(true) {} // NOLINT @@ -163,16 +162,13 @@ class ParserImpl { /// @param e the Expect to copy this Maybe from template inline Maybe(const Expect& e) // NOLINT - : value(e.value), source(e.value), errored(e.errored), matched(!e.errored) {} + : value(e.value), errored(e.errored), matched(!e.errored) {} /// Move from an Expect. /// @param e the Expect to move this Maybe from template inline Maybe(Expect&& e) // NOLINT - : value(std::move(e.value)), - source(std::move(e.source)), - errored(e.errored), - matched(!e.errored) {} + : value(std::move(e.value)), errored(e.errored), matched(!e.errored) {} /// Copy constructor inline Maybe(const Maybe&) = default; @@ -197,8 +193,6 @@ class ParserImpl { /// The value of a successful parse. /// Zero-initialized when there was a parse error. T value{}; - /// Optional source of the value. - Source source; /// True if there was a error parsing. bool errored = false; /// True if there was a error parsing. @@ -268,33 +262,14 @@ class ParserImpl { /// VarDeclInfo contains the parsed information for variable declaration. struct VarDeclInfo { - /// Constructor - VarDeclInfo(); - /// Copy constructor - /// @param other the VarDeclInfo to copy - VarDeclInfo(const VarDeclInfo& other); - /// Constructor - /// @param source_in variable declaration source - /// @param name_in variable name - /// @param address_space_in variable address space - /// @param access_in variable access control - /// @param type_in variable type - VarDeclInfo(Source source_in, - std::string name_in, - type::AddressSpace address_space_in, - type::Access access_in, - ast::Type type_in); - /// Destructor - ~VarDeclInfo(); - /// Variable declaration source Source source; /// Variable name std::string name; /// Variable address space - type::AddressSpace address_space = type::AddressSpace::kUndefined; + const ast::Expression* address_space = nullptr; /// Variable access control - type::Access access = type::Access::kUndefined; + const ast::Expression* access = nullptr; /// Variable type ast::Type type; }; @@ -302,9 +277,9 @@ class ParserImpl { /// VariableQualifier contains the parsed information for a variable qualifier struct VariableQualifier { /// The variable's address space - type::AddressSpace address_space = type::AddressSpace::kUndefined; + const ast::Expression* address_space = nullptr; /// The variable's access control - type::Access access = type::Access::kUndefined; + const ast::Expression* access = nullptr; }; /// MatrixDimensions contains the column and row information for a matrix @@ -447,10 +422,6 @@ class ParserImpl { /// Parses a `type_specifier` grammar element /// @returns the parsed Type or nullptr if none matched. Maybe type_specifier(); - /// Parses an `address_space` grammar element, erroring on parse failure. - /// @param use a description of what was being parsed if an error was raised. - /// @returns the address space or type::AddressSpace::kUndefined if none matched - Expect expect_address_space(std::string_view use); /// Parses a `struct_decl` grammar element. /// @returns the struct type or nullptr on error Maybe struct_decl(); @@ -482,11 +453,6 @@ class ParserImpl { /// not match a stage name. /// @returns the pipeline stage. Expect expect_pipeline_stage(); - /// Parses an access control identifier, erroring if the next token does not - /// match a valid access control. - /// @param use a description of what was being parsed if an error was raised - /// @returns the parsed access control. - Expect expect_access_mode(std::string_view use); /// Parses an interpolation sample name identifier, erroring if the next token does not match a /// valid sample name. /// @returns the parsed sample name. @@ -597,8 +563,9 @@ class ParserImpl { /// @returns the parsed expression or nullptr Maybe expression(); /// Parses the `expression` grammar rule + /// @param use the use of the expression /// @returns the parsed expression or error - Expect expect_expression(); + Expect expect_expression(std::string_view use); /// Parses a comma separated expression list /// @param use the use of the expression list /// @param terminator the terminating token for the list diff --git a/src/tint/reader/wgsl/parser_impl_address_space_test.cc b/src/tint/reader/wgsl/parser_impl_address_space_test.cc deleted file mode 100644 index 32c3c47c69..0000000000 --- a/src/tint/reader/wgsl/parser_impl_address_space_test.cc +++ /dev/null @@ -1,62 +0,0 @@ -// 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/reader/wgsl/parser_impl_test_helper.h" - -namespace tint::reader::wgsl { -namespace { - -struct AddressSpaceData { - const char* input; - type::AddressSpace result; -}; -inline std::ostream& operator<<(std::ostream& out, AddressSpaceData data) { - out << std::string(data.input); - return out; -} - -class ParserAddressSpaceTest : public ParserImplTestWithParam {}; - -TEST_P(ParserAddressSpaceTest, Parses) { - auto params = GetParam(); - auto p = parser(params.input); - - auto sc = p->expect_address_space("test"); - EXPECT_FALSE(sc.errored); - EXPECT_FALSE(p->has_error()); - EXPECT_EQ(sc.value, params.result); - - auto& t = p->next(); - EXPECT_TRUE(t.IsEof()); -} -INSTANTIATE_TEST_SUITE_P( - ParserImplTest, - ParserAddressSpaceTest, - testing::Values(AddressSpaceData{"uniform", type::AddressSpace::kUniform}, - AddressSpaceData{"workgroup", type::AddressSpace::kWorkgroup}, - AddressSpaceData{"storage", type::AddressSpace::kStorage}, - AddressSpaceData{"private", type::AddressSpace::kPrivate}, - AddressSpaceData{"function", type::AddressSpace::kFunction})); - -TEST_F(ParserImplTest, AddressSpace_NoMatch) { - auto p = parser("not-a-address-space"); - auto sc = p->expect_address_space("test"); - EXPECT_EQ(sc.errored, true); - EXPECT_TRUE(p->has_error()); - EXPECT_EQ(p->error(), R"(1:1: expected address space for test -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')"); -} - -} // namespace -} // namespace tint::reader::wgsl diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc index d40898aede..18b40beb8f 100644 --- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc +++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc @@ -799,7 +799,7 @@ var x : texture_storage_2d;", - R"(test.wgsl:1:28 error: expected expression + R"(test.wgsl:1:28 error: expected expression for type template argument list var x : texture_storage_2d<>; ^ )"); @@ -1123,20 +1123,11 @@ var i : ptr i : i32", - R"(test.wgsl:1:5 error: expected address space for variable declaration -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup' -var i : i32 - ^^^^ -)"); -} - TEST_F(ParserImplErrorTest, GlobalDeclVarStorageDeclMissingGThan) { EXPECT("var' for variable declaration + R"(test.wgsl:1:4 error: missing closing '>' for variable declaration varbuilder().Symbols(), var->name, "a"); - ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32"); - - EXPECT_EQ(var->declared_address_space, type::AddressSpace::kPrivate); + ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "private"); EXPECT_EQ(var->source.range.begin.line, 1u); EXPECT_EQ(var->source.range.begin.column, 14u); @@ -58,8 +56,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithInitializer) { ast::CheckIdentifier(p->builder().Symbols(), var->name, "a"); ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32"); - - EXPECT_EQ(var->declared_address_space, type::AddressSpace::kPrivate); + ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "private"); EXPECT_EQ(var->source.range.begin.line, 1u); EXPECT_EQ(var->source.range.begin.column, 14u); @@ -83,11 +80,8 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithAttribute) { ASSERT_NE(var, nullptr); ast::CheckIdentifier(p->builder().Symbols(), var->name, "a"); - ASSERT_NE(var->type, nullptr); - ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32"); - - EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform); + ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform"); EXPECT_EQ(var->source.range.begin.line, 1u); EXPECT_EQ(var->source.range.begin.column, 36u); @@ -116,10 +110,8 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithAttribute_MulitpleGroups) { ASSERT_NE(var, nullptr); ast::CheckIdentifier(p->builder().Symbols(), var->name, "a"); - ASSERT_NE(var->type, nullptr); ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32"); - - EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform); + ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform"); EXPECT_EQ(var->source.range.begin.line, 1u); EXPECT_EQ(var->source.range.begin.column, 36u); @@ -162,19 +154,5 @@ TEST_F(ParserImplTest, GlobalVariableDecl_InvalidConstExpr) { EXPECT_EQ(p->error(), "1:24: missing initializer for 'var' declaration"); } -TEST_F(ParserImplTest, GlobalVariableDecl_InvalidVariableDecl) { - auto p = parser("var a : f32;"); - auto attrs = p->attribute_list(); - EXPECT_FALSE(attrs.errored); - EXPECT_FALSE(attrs.matched); - auto e = p->global_variable_decl(attrs.value); - EXPECT_TRUE(p->has_error()); - EXPECT_TRUE(e.errored); - EXPECT_FALSE(e.matched); - EXPECT_EQ(e.value, nullptr); - EXPECT_EQ(p->error(), R"(1:5: expected address space for variable declaration -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')"); -} - } // namespace } // namespace tint::reader::wgsl diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc index b40d887a4c..25e60e0368 100644 --- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc +++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc @@ -137,7 +137,7 @@ TEST_P(VecMissingType, Handles_Missing_Type) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), "1:6: expected expression"); + ASSERT_EQ(p->error(), "1:6: expected expression for type template argument list"); } INSTANTIATE_TEST_SUITE_P(ParserImplTest, VecMissingType, @@ -211,7 +211,7 @@ TEST_F(ParserImplTest, TypeDecl_Ptr_MissingAddressSpace) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), R"(1:5: expected expression)"); + ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)"); } TEST_F(ParserImplTest, TypeDecl_Ptr_MissingParams) { @@ -221,7 +221,7 @@ TEST_F(ParserImplTest, TypeDecl_Ptr_MissingParams) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), R"(1:5: expected expression)"); + ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)"); } TEST_F(ParserImplTest, TypeDecl_Atomic) { @@ -256,7 +256,7 @@ TEST_F(ParserImplTest, TypeDecl_Atomic_MissingType) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), "1:8: expected expression"); + ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list"); } TEST_F(ParserImplTest, TypeDecl_Array_AbstractIntLiteralSize) { @@ -431,7 +431,7 @@ TEST_P(MatrixMissingType, Handles_Missing_Type) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), "1:8: expected expression"); + ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list"); } INSTANTIATE_TEST_SUITE_P(ParserImplTest, MatrixMissingType, diff --git a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc index 9138d55a7a..993ecb0ad9 100644 --- a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc +++ b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc @@ -84,8 +84,7 @@ TEST_F(ParserImplTest, VariableDecl_WithAddressSpace) { EXPECT_EQ(v->name, "my_var"); ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32"); - - EXPECT_EQ(v->address_space, type::AddressSpace::kPrivate); + ast::CheckIdentifier(p->builder().Symbols(), v->address_space, "private"); EXPECT_EQ(v->source.range.begin.line, 1u); EXPECT_EQ(v->source.range.begin.column, 14u); @@ -102,20 +101,7 @@ TEST_F(ParserImplTest, VariableDecl_WithPushConstant) { EXPECT_EQ(v->name, "my_var"); ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32"); - - EXPECT_EQ(v->address_space, type::AddressSpace::kPushConstant); -} - -TEST_F(ParserImplTest, VariableDecl_InvalidAddressSpace) { - auto p = parser("var my_var : f32"); - auto v = p->variable_decl(); - EXPECT_FALSE(v.matched); - EXPECT_TRUE(v.errored); - EXPECT_TRUE(p->has_error()); - EXPECT_EQ(p->error(), - R"(1:5: expected address space for variable declaration -Did you mean 'uniform'? -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')"); + ast::CheckIdentifier(p->builder().Symbols(), v->address_space, "push_constant"); } } // namespace diff --git a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc index 111f20787a..9a8b936e67 100644 --- a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc +++ b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc @@ -12,6 +12,7 @@ // See the License for the specific language governing permissions and // limitations under the License. +#include "src/tint/ast/test_helper.h" #include "src/tint/reader/wgsl/parser_impl_test_helper.h" namespace tint::reader::wgsl { @@ -31,14 +32,23 @@ class VariableQualifierTest : public ParserImplTestWithParam"); + auto p = parser(std::string("var<") + params.input + "> name"); - auto sc = p->variable_qualifier(); + auto sc = p->variable_decl(); EXPECT_FALSE(p->has_error()); EXPECT_FALSE(sc.errored); EXPECT_TRUE(sc.matched); - EXPECT_EQ(sc->address_space, params.address_space); - EXPECT_EQ(sc->access, params.access); + if (params.address_space != type::AddressSpace::kUndefined) { + ast::CheckIdentifier(p->builder().Symbols(), sc->address_space, + utils::ToString(params.address_space)); + } else { + EXPECT_EQ(sc->address_space, nullptr); + } + if (params.access != type::Access::kUndefined) { + ast::CheckIdentifier(p->builder().Symbols(), sc->access, utils::ToString(params.access)); + } else { + EXPECT_EQ(sc->access, nullptr); + } auto& t = p->next(); EXPECT_TRUE(t.IsEof()); @@ -57,24 +67,13 @@ INSTANTIATE_TEST_SUITE_P( VariableStorageData{"storage, read_write", type::AddressSpace::kStorage, type::Access::kReadWrite})); -TEST_F(ParserImplTest, VariableQualifier_NoMatch) { - auto p = parser(""); - auto sc = p->variable_qualifier(); - EXPECT_TRUE(p->has_error()); - EXPECT_TRUE(sc.errored); - EXPECT_FALSE(sc.matched); - EXPECT_EQ(p->error(), R"(1:2: expected address space for variable declaration -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')"); -} - TEST_F(ParserImplTest, VariableQualifier_Empty) { - auto p = parser("<>"); - auto sc = p->variable_qualifier(); + auto p = parser("var<> name"); + auto sc = p->variable_decl(); EXPECT_TRUE(p->has_error()); EXPECT_TRUE(sc.errored); EXPECT_FALSE(sc.matched); - EXPECT_EQ(p->error(), R"(1:2: expected address space for variable declaration -Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')"); + EXPECT_EQ(p->error(), R"(1:5: expected expression for 'var' address space)"); } TEST_F(ParserImplTest, VariableQualifier_MissingLessThan) { @@ -105,7 +104,7 @@ TEST_F(ParserImplTest, VariableQualifier_MissingGreaterThan) { EXPECT_TRUE(p->has_error()); EXPECT_TRUE(sc.errored); EXPECT_FALSE(sc.matched); - EXPECT_EQ(p->error(), "1:9: expected '>' for variable declaration"); + EXPECT_EQ(p->error(), "1:1: missing closing '>' for variable declaration"); } } // namespace diff --git a/src/tint/resolver/address_space_validation_test.cc b/src/tint/resolver/address_space_validation_test.cc index 2c879de526..95301c74ad 100644 --- a/src/tint/resolver/address_space_validation_test.cc +++ b/src/tint/resolver/address_space_validation_test.cc @@ -32,8 +32,9 @@ TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_NoAddressSpace_Fail) { GlobalVar(Source{{12, 34}}, "g", ty.f32()); EXPECT_FALSE(r()->Resolve()); - EXPECT_EQ(r()->error(), - "12:34 error: module-scope 'var' declaration must have a address space"); + EXPECT_EQ( + r()->error(), + R"(12:34 error: module-scope 'var' declarations that are not of texture or sampler types must provide an address space)"); } TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NoAddressSpace_Fail) { @@ -469,7 +470,7 @@ TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_NotStorage_AccessMode) EXPECT_EQ( r()->error(), - R"(12:34 error: only variables in address space may declare an access mode)"); + R"(12:34 error: only variables in address space may specify an access mode)"); } TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NotStorage_AccessMode) { @@ -481,7 +482,7 @@ TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NotStorage_AccessMode) { EXPECT_EQ( r()->error(), - R"(12:34 error: only pointers in address space may declare an access mode)"); + R"(12:34 error: only pointers in address space may specify an access mode)"); } TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_Storage_ReadAccessMode) { diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc index 16a3f8ec15..75e377fcab 100644 --- a/src/tint/resolver/dependency_graph.cc +++ b/src/tint/resolver/dependency_graph.cc @@ -183,13 +183,9 @@ class DependencyScanner { Declare(func->name->symbol, func); TraverseFunction(func); }, - [&](const ast::Variable* var) { - Declare(var->name->symbol, var); - TraverseTypeExpression(var->type); - TraverseAttributes(var->attributes); - if (var->initializer) { - TraverseValueExpression(var->initializer); - } + [&](const ast::Variable* v) { + Declare(v->name->symbol, v); + TraverseVariable(v); }, [&](const ast::DiagnosticDirective*) { // Diagnostic directives do not affect the dependency graph. @@ -204,8 +200,18 @@ class DependencyScanner { } private: - /// Traverses the function, performing symbol resolution and determining - /// global dependencies. + /// Traverses the variable, performing symbol resolution. + void TraverseVariable(const ast::Variable* v) { + if (auto* var = v->As()) { + TraverseAddressSpaceExpression(var->declared_address_space); + TraverseAccessExpression(var->declared_access); + } + TraverseTypeExpression(v->type); + TraverseAttributes(v->attributes); + TraverseValueExpression(v->initializer); + } + + /// Traverses the function, performing symbol resolution and determining global dependencies. void TraverseFunction(const ast::Function* func) { TraverseAttributes(func->attributes); TraverseAttributes(func->return_type_attributes); @@ -301,8 +307,7 @@ class DependencyScanner { if (auto* shadows = scope_stack_.Get(v->variable->name->symbol)) { graph_.shadows.Add(v->variable, shadows); } - TraverseTypeExpression(v->variable->type); - TraverseValueExpression(v->variable->initializer); + TraverseVariable(v->variable); Declare(v->variable->name->symbol, v->variable); }, [&](const ast::WhileStatement* w) { @@ -345,6 +350,18 @@ class DependencyScanner { TraverseExpression(root, "type", "references"); } + /// Traverses the expression @p root_expr for the intended use as an address space, performing + /// symbol resolution and determining global dependencies. + void TraverseAddressSpaceExpression(const ast::Expression* root) { + TraverseExpression(root, "address space", "references"); + } + + /// Traverses the expression @p root_expr for the intended use as an access, performing symbol + /// resolution and determining global dependencies. + void TraverseAccessExpression(const ast::Expression* root) { + TraverseExpression(root, "access", "references"); + } + /// Traverses the expression @p root_expr for the intended use as a call target, performing /// symbol resolution and determining global dependencies. void TraverseCallableExpression(const ast::Expression* root) { diff --git a/src/tint/resolver/dependency_graph_test.cc b/src/tint/resolver/dependency_graph_test.cc index 31e484eec4..a316ad0443 100644 --- a/src/tint/resolver/dependency_graph_test.cc +++ b/src/tint/resolver/dependency_graph_test.cc @@ -1437,13 +1437,13 @@ TEST_P(ResolverDependencyGraphResolveToAddressSpace, Resolve) { << resolved->String(Symbols(), Diagnostics()); } -TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalVar) { +TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalConst) { const auto use = std::get<0>(GetParam()); const auto builtin = std::get<1>(GetParam()); const auto symbol = Symbols().New(utils::ToString(builtin)); SymbolTestHelper helper(this); - auto* decl = helper.Add(SymbolDeclKind::GlobalVar, symbol); + auto* decl = helper.Add(SymbolDeclKind::GlobalConst, symbol); auto* ident = helper.Add(use, symbol); helper.Build(); diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc index b04831ee62..1062958a7d 100644 --- a/src/tint/resolver/function_validation_test.cc +++ b/src/tint/resolver/function_validation_test.cc @@ -1088,15 +1088,14 @@ TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceWithExtension) { param.expectation == Expectation::kPassWithFullPtrParameterExtension) { ASSERT_TRUE(r()->Resolve()) << r()->error(); } else { - std::stringstream ss; - ss << param.address_space; EXPECT_FALSE(r()->Resolve()); if (param.expectation == Expectation::kInvalid) { - EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" + ss.str() + "'"); + EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" + + utils::ToString(param.address_space) + "'"); } else { EXPECT_EQ(r()->error(), - "12:34 error: function parameter of pointer type cannot be in '" + ss.str() + - "' address space"); + "12:34 error: function parameter of pointer type cannot be in '" + + utils::ToString(param.address_space) + "' address space"); } } } @@ -1105,8 +1104,8 @@ INSTANTIATE_TEST_SUITE_P( ResolverFunctionParameterValidationTest, testing::Values( TestParams{type::AddressSpace::kUndefined, Expectation::kInvalid}, - TestParams{type::AddressSpace::kIn, Expectation::kInvalid}, - TestParams{type::AddressSpace::kOut, Expectation::kInvalid}, + TestParams{type::AddressSpace::kIn, Expectation::kAlwaysFail}, + TestParams{type::AddressSpace::kOut, Expectation::kAlwaysFail}, TestParams{type::AddressSpace::kUniform, Expectation::kPassWithFullPtrParameterExtension}, TestParams{type::AddressSpace::kWorkgroup, Expectation::kPassWithFullPtrParameterExtension}, TestParams{type::AddressSpace::kHandle, Expectation::kInvalid}, diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc index b87852f646..833676577a 100644 --- a/src/tint/resolver/resolver.cc +++ b/src/tint/resolver/resolver.cc @@ -473,8 +473,14 @@ sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) { return nullptr; } - auto address_space = var->declared_address_space; - if (address_space == type::AddressSpace::kUndefined) { + auto address_space = type::AddressSpace::kUndefined; + if (var->declared_address_space) { + auto expr = AddressSpaceExpression(var->declared_address_space); + if (!expr) { + return nullptr; + } + address_space = expr->Value(); + } else { // No declared address space. Infer from usage / type. if (!is_global) { address_space = type::AddressSpace::kFunction; @@ -494,8 +500,14 @@ sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) { return nullptr; } - auto access = var->declared_access; - if (access == type::Access::kUndefined) { + auto access = type::Access::kUndefined; + if (var->declared_access) { + auto expr = AccessExpression(var->declared_access); + if (!expr) { + return nullptr; + } + access = expr->Value(); + } else { access = DefaultAccessForAddressSpace(address_space); } diff --git a/src/tint/resolver/validation_test.cc b/src/tint/resolver/validation_test.cc index 73fa9b879e..8ee795ce91 100644 --- a/src/tint/resolver/validation_test.cc +++ b/src/tint/resolver/validation_test.cc @@ -295,22 +295,23 @@ TEST_F(ResolverValidationTest, AddressSpace_FunctionVariableI32) { TEST_F(ResolverValidationTest, AddressSpace_SamplerExplicitAddressSpace) { auto t = ty.sampler(type::SamplerKind::kSampler); - GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a)); + GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kPrivate, Binding(0_a), Group(0_a)); EXPECT_FALSE(r()->Resolve()); EXPECT_EQ(r()->error(), - R"(12:34 error: variables of type 'sampler' must not have a address space)"); + R"(12:34 error: variables of type 'sampler' must not specifiy an address space)"); } TEST_F(ResolverValidationTest, AddressSpace_TextureExplicitAddressSpace) { auto t = ty.sampled_texture(type::TextureDimension::k1d, ty.f32()); - GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a)); + GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kFunction, Binding(0_a), Group(0_a)); EXPECT_FALSE(r()->Resolve()) << r()->error(); - EXPECT_EQ(r()->error(), - R"(12:34 error: variables of type 'texture_1d' must not have a address space)"); + EXPECT_EQ( + r()->error(), + R"(12:34 error: variables of type 'texture_1d' must not specifiy an address space)"); } TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) { diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index 5b399cfb14..2935c70c24 100644 --- a/src/tint/resolver/validator.cc +++ b/src/tint/resolver/validator.cc @@ -299,7 +299,7 @@ bool Validator::Pointer(const ast::TemplatedIdentifier* a, const type::Pointer* // * For the storage address space, the access mode is optional, and defaults to read. // * For other address spaces, the access mode must not be written. if (s->AddressSpace() != type::AddressSpace::kStorage) { - AddError("only pointers in address space may declare an access mode", + AddError("only pointers in address space may specify an access mode", a->source); return false; } @@ -604,7 +604,7 @@ bool Validator::GlobalVariable( } bool ok = Switch( decl, // - [&](const ast::Var*) { + [&](const ast::Var* var) { if (auto* init = global->Initializer(); init && init->Stage() > sem::EvaluationStage::kOverride) { AddError("module-scope 'var' initializer must be a constant or override-expression", @@ -612,8 +612,11 @@ bool Validator::GlobalVariable( return false; } - if (global->AddressSpace() == type::AddressSpace::kUndefined) { - AddError("module-scope 'var' declaration must have a address space", decl->source); + if (!var->declared_address_space && !global->Type()->UnwrapRef()->is_handle()) { + AddError( + "module-scope 'var' declarations that are not of texture or sampler types must " + "provide an address space", + decl->source); return false; } @@ -696,25 +699,23 @@ bool Validator::Var(const sem::Variable* v) const { return false; } - if (store_ty->is_handle()) { - if (var->declared_address_space != type::AddressSpace::kUndefined) { - // 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. - AddError("variables of type '" + sem_.TypeNameOf(store_ty) + - "' must not have a address space", - var->source); - return false; - } + if (store_ty->is_handle() && var->declared_address_space) { + // 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. + AddError("variables of type '" + sem_.TypeNameOf(store_ty) + + "' must not specifiy an address space", + var->source); + return false; } - if (var->declared_access != type::Access::kUndefined) { + if (var->declared_access) { // https://www.w3.org/TR/WGSL/#access-mode-defaults // When writing a variable declaration or a pointer type in WGSL source: // * For the storage address space, the access mode is optional, and defaults to read. // * For other address spaces, the access mode must not be written. - if (var->declared_address_space != type::AddressSpace::kStorage) { - AddError("only variables in address space may declare an access mode", + if (v->AddressSpace() != type::AddressSpace::kStorage) { + AddError("only variables in address space may specify an access mode", var->source); return false; } @@ -726,8 +727,8 @@ bool Validator::Var(const sem::Variable* v) const { } if (IsValidationEnabled(var->attributes, ast::DisabledValidation::kIgnoreAddressSpace) && - (var->declared_address_space == type::AddressSpace::kIn || - var->declared_address_space == type::AddressSpace::kOut)) { + (v->AddressSpace() == type::AddressSpace::kIn || + v->AddressSpace() == type::AddressSpace::kOut)) { AddError("invalid use of input/output address space", var->source); return false; } diff --git a/src/tint/transform/binding_remapper.cc b/src/tint/transform/binding_remapper.cc index b6579807e3..192b2c93fd 100644 --- a/src/tint/transform/binding_remapper.cc +++ b/src/tint/transform/binding_remapper.cc @@ -122,11 +122,12 @@ Transform::ApplyResult BindingRemapper::Apply(const Program* src, // Replace any access controls. auto ac_it = remappings->access_controls.find(from); if (ac_it != remappings->access_controls.end()) { - type::Access ac = ac_it->second; - if (ac == type::Access::kUndefined) { - b.Diagnostics().add_error( - diag::System::Transform, - "invalid access mode (" + std::to_string(static_cast(ac)) + ")"); + type::Access access = ac_it->second; + if (access == type::Access::kUndefined) { + b.Diagnostics().add_error(diag::System::Transform, + "invalid access mode (" + + std::to_string(static_cast(access)) + + ")"); return Program(std::move(b)); } auto* sem = src->Sem().Get(var); @@ -139,9 +140,14 @@ Transform::ApplyResult BindingRemapper::Apply(const Program* src, } auto* ty = sem->Type()->UnwrapRef(); auto inner_ty = CreateASTTypeFor(ctx, ty); - auto* new_var = b.Var(ctx.Clone(var->source), ctx.Clone(var->name->symbol), - inner_ty, var->declared_address_space, ac, - ctx.Clone(var->initializer), ctx.Clone(var->attributes)); + auto* new_var = + b.create(ctx.Clone(var->source), // source + b.Ident(ctx.Clone(var->name->symbol)), // name + inner_ty, // type + ctx.Clone(var->declared_address_space), // address space + b.Expr(access), // access + ctx.Clone(var->initializer), // initializer + ctx.Clone(var->attributes)); // attributes ctx.Replace(var, new_var); } diff --git a/src/tint/transform/canonicalize_entry_point_io_test.cc b/src/tint/transform/canonicalize_entry_point_io_test.cc index 5af8a142fd..f0ce8c34c5 100644 --- a/src/tint/transform/canonicalize_entry_point_io_test.cc +++ b/src/tint/transform/canonicalize_entry_point_io_test.cc @@ -67,11 +67,11 @@ fn frag_main(@location(1) loc1 : f32, )"; auto* expect = R"( -@location(1) @internal(disable_validation__ignore_address_space) var loc1_1 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var loc2_1 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4; -@builtin(position) @internal(disable_validation__ignore_address_space) var coord_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4; fn frag_main_inner(loc1 : f32, loc2 : vec4, coord : vec4) { var col : f32 = (coord.x * loc1); @@ -251,13 +251,13 @@ fn frag_main(@location(0) loc0 : f32, )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var loc0_1 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> loc0_1 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var loc1_1 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var loc2_1 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4; -@builtin(position) @internal(disable_validation__ignore_address_space) var coord_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4; struct FragBuiltins { coord : vec4, @@ -304,13 +304,13 @@ struct FragLocations { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var loc0_1 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> loc0_1 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var loc1_1 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var loc2_1 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4; -@builtin(position) @internal(disable_validation__ignore_address_space) var coord_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4; fn frag_main_inner(loc0 : f32, locations : FragLocations, builtins : FragBuiltins) { var col : f32 = ((builtins.coord.x * locations.loc1) + loc0); @@ -567,7 +567,7 @@ fn frag_main() -> @builtin(frag_depth) f32 { )"; auto* expect = R"( -@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var value : f32; +@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> value : f32; fn frag_main_inner() -> f32 { return 1.0; @@ -674,11 +674,11 @@ fn frag_main() -> FragOutput { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var color_1 : vec4; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4; -@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var depth_1 : f32; +@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> depth_1 : f32; -@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var mask_1 : array; +@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array; struct FragOutput { color : vec4, @@ -729,11 +729,11 @@ struct FragOutput { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var color_1 : vec4; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4; -@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var depth_1 : f32; +@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> depth_1 : f32; -@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var mask_1 : array; +@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array; fn frag_main_inner() -> FragOutput { var output : FragOutput; @@ -1028,13 +1028,13 @@ fn frag_main2(inputs : FragmentInput) { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var value_1 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_1 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var mul_1 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_1 : f32; -@location(0) @internal(disable_validation__ignore_address_space) var value_2 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_2 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var mul_2 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_2 : f32; struct FragmentInput { value : f32, @@ -1094,13 +1094,13 @@ struct FragmentInput { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var value_1 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_1 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var mul_1 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_1 : f32; -@location(0) @internal(disable_validation__ignore_address_space) var value_2 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_2 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var mul_2 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_2 : f32; fn frag_main1_inner(inputs : FragmentInput) { var x : f32 = foo(inputs); @@ -1952,39 +1952,39 @@ fn frag_main(inputs : FragmentInterface) -> FragmentInterface { auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var i_1 : i32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> i_1 : i32; -@location(1) @internal(disable_validation__ignore_address_space) var u_1 : u32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> u_1 : u32; -@location(2) @internal(disable_validation__ignore_address_space) var vi_1 : vec4; +@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4; -@location(3) @internal(disable_validation__ignore_address_space) var vu_1 : vec4; +@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4; -@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var i_2 : i32; +@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> i_2 : i32; -@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var u_2 : u32; +@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> u_2 : u32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vi_2 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4; -@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vu_2 : vec4; +@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4; -@builtin(position) @internal(disable_validation__ignore_address_space) var pos_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4; -@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var i_3 : i32; +@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> i_3 : i32; -@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var u_3 : u32; +@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> u_3 : u32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vi_3 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4; -@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vu_3 : vec4; +@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4; -@location(0) @internal(disable_validation__ignore_address_space) var i_4 : i32; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> i_4 : i32; -@location(1) @internal(disable_validation__ignore_address_space) var u_4 : u32; +@location(1) @internal(disable_validation__ignore_address_space) var<__out> u_4 : u32; -@location(2) @internal(disable_validation__ignore_address_space) var vi_4 : vec4; +@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4; -@location(3) @internal(disable_validation__ignore_address_space) var vu_4 : vec4; +@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4; struct VertexIn { i : i32, @@ -2082,39 +2082,39 @@ struct FragmentInterface { auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var i_1 : i32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> i_1 : i32; -@location(1) @internal(disable_validation__ignore_address_space) var u_1 : u32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> u_1 : u32; -@location(2) @internal(disable_validation__ignore_address_space) var vi_1 : vec4; +@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4; -@location(3) @internal(disable_validation__ignore_address_space) var vu_1 : vec4; +@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4; -@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var i_2 : i32; +@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> i_2 : i32; -@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var u_2 : u32; +@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> u_2 : u32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vi_2 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4; -@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vu_2 : vec4; +@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4; -@builtin(position) @internal(disable_validation__ignore_address_space) var pos_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4; -@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var i_3 : i32; +@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> i_3 : i32; -@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var u_3 : u32; +@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> u_3 : u32; -@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vi_3 : vec4; +@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4; -@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var vu_3 : vec4; +@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4; -@location(0) @internal(disable_validation__ignore_address_space) var i_4 : i32; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> i_4 : i32; -@location(1) @internal(disable_validation__ignore_address_space) var u_4 : u32; +@location(1) @internal(disable_validation__ignore_address_space) var<__out> u_4 : u32; -@location(2) @internal(disable_validation__ignore_address_space) var vi_4 : vec4; +@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4; -@location(3) @internal(disable_validation__ignore_address_space) var vu_4 : vec4; +@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4; fn vert_main_inner(in : VertexIn) -> VertexOut { return VertexOut(in.i, in.u, in.vi, in.vu, vec4()); @@ -3161,9 +3161,9 @@ fn vert_main() -> @builtin(position) vec4 { )"; auto* expect = R"( -@builtin(position) @internal(disable_validation__ignore_address_space) var value : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> value : vec4; -@builtin(point_size) @internal(disable_validation__ignore_address_space) var vertex_point_size : f32; +@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32; fn vert_main_inner() -> vec4 { return vec4(); @@ -3236,9 +3236,9 @@ fn vert_main() -> VertOut { )"; auto* expect = R"( -@builtin(position) @internal(disable_validation__ignore_address_space) var pos_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4; -@builtin(point_size) @internal(disable_validation__ignore_address_space) var vertex_point_size : f32; +@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32; struct VertOut { pos : vec4, @@ -3277,9 +3277,9 @@ struct VertOut { )"; auto* expect = R"( -@builtin(position) @internal(disable_validation__ignore_address_space) var pos_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4; -@builtin(point_size) @internal(disable_validation__ignore_address_space) var vertex_point_size : f32; +@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32; fn vert_main_inner() -> VertOut { return VertOut(); @@ -3424,15 +3424,15 @@ fn vert_main(collide : VertIn1, collide_1 : VertIn2) -> VertOut { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var collide_2 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> collide_2 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var collide_3 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> collide_3 : f32; -@location(0) @internal(disable_validation__ignore_address_space) var vertex_point_size_3 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_3 : f32; -@builtin(position) @internal(disable_validation__ignore_address_space) var vertex_point_size_1_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4; -@builtin(point_size) @internal(disable_validation__ignore_address_space) var vertex_point_size_4 : f32; +@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_4 : f32; var vertex_point_size : f32; @@ -3502,15 +3502,15 @@ struct VertOut { )"; auto* expect = R"( -@location(0) @internal(disable_validation__ignore_address_space) var collide_2 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__in> collide_2 : f32; -@location(1) @internal(disable_validation__ignore_address_space) var collide_3 : f32; +@location(1) @internal(disable_validation__ignore_address_space) var<__in> collide_3 : f32; -@location(0) @internal(disable_validation__ignore_address_space) var vertex_point_size_3 : f32; +@location(0) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_3 : f32; -@builtin(position) @internal(disable_validation__ignore_address_space) var vertex_point_size_1_1 : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4; -@builtin(point_size) @internal(disable_validation__ignore_address_space) var vertex_point_size_4 : f32; +@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_4 : f32; fn vert_main_inner(collide : VertIn1, collide_1 : VertIn2) -> VertOut { let x = (collide.collide + collide_1.collide); @@ -3868,11 +3868,11 @@ fn main(@builtin(sample_index) sample_index : u32, )"; auto* expect = R"( -@builtin(sample_index) @interpolate(flat) @internal(disable_validation__ignore_address_space) var sample_index_1 : u32; +@builtin(sample_index) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> sample_index_1 : u32; -@builtin(sample_mask) @interpolate(flat) @internal(disable_validation__ignore_address_space) var mask_in_1 : array; +@builtin(sample_mask) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> mask_in_1 : array; -@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var value : array; +@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> value : array; fn main_inner(sample_index : u32, mask_in : u32) -> u32 { return mask_in; @@ -3903,11 +3903,11 @@ fn fragment_main(@builtin(sample_index) sample_index : u32, )"; auto* expect = R"( -@builtin(sample_index) @internal(disable_validation__ignore_address_space) var gl_SampleID : i32; +@builtin(sample_index) @internal(disable_validation__ignore_address_space) var<__in> gl_SampleID : i32; -@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var gl_SampleMaskIn : array; +@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__in> gl_SampleMaskIn : array; -@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var gl_SampleMask : array; +@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> gl_SampleMask : array; fn fragment_main(sample_index : u32, mask_in : u32) -> u32 { return mask_in; @@ -3938,11 +3938,11 @@ fn vertex_main(@builtin(vertex_index) vertexID : u32, )"; auto* expect = R"( -@builtin(vertex_index) @internal(disable_validation__ignore_address_space) var gl_VertexID : i32; +@builtin(vertex_index) @internal(disable_validation__ignore_address_space) var<__in> gl_VertexID : i32; -@builtin(instance_index) @internal(disable_validation__ignore_address_space) var gl_InstanceID : i32; +@builtin(instance_index) @internal(disable_validation__ignore_address_space) var<__in> gl_InstanceID : i32; -@builtin(position) @internal(disable_validation__ignore_address_space) var gl_Position : vec4; +@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> gl_Position : vec4; fn vertex_main(vertexID : u32, instanceID : u32) -> vec4 { return vec4((f32(vertexID) + f32(instanceID))); diff --git a/src/tint/transform/clamp_frag_depth.cc b/src/tint/transform/clamp_frag_depth.cc index 57d9a40a3c..8d84ed3654 100644 --- a/src/tint/transform/clamp_frag_depth.cc +++ b/src/tint/transform/clamp_frag_depth.cc @@ -88,10 +88,11 @@ Transform::ApplyResult ClampFragDepth::Apply(const Program* src, const DataMap&, // Abort on any use of push constants in the module. for (auto* global : src->AST().GlobalVariables()) { if (auto* var = global->As()) { - if (TINT_UNLIKELY(var->declared_address_space == type::AddressSpace::kPushConstant)) { + auto* v = src->Sem().Get(var); + if (TINT_UNLIKELY(v->AddressSpace() == type::AddressSpace::kPushConstant)) { TINT_ICE(Transform, b.Diagnostics()) << "ClampFragDepth doesn't know how to handle module that already use push " - "constants."; + "constants"; return Program(std::move(b)); } } diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.cc b/src/tint/transform/module_scope_var_to_entry_point_param.cc index ff65fecf1e..47356e9626 100644 --- a/src/tint/transform/module_scope_var_to_entry_point_param.cc +++ b/src/tint/transform/module_scope_var_to_entry_point_param.cc @@ -117,7 +117,6 @@ struct ModuleScopeVarToEntryPointParam::State { WorkgroupParameterMemberList& workgroup_parameter_members, bool& is_pointer, bool& is_wrapped) { - auto* var_ast = var->Declaration()->As(); auto* ty = var->Type()->UnwrapRef(); // Helper to create an AST node for the store type of the variable. @@ -160,7 +159,9 @@ struct ModuleScopeVarToEntryPointParam::State { is_wrapped = true; } - param_type = ctx.dst->ty.pointer(param_type, sc, var_ast->declared_access); + param_type = sc == type::AddressSpace::kStorage + ? ctx.dst->ty.pointer(param_type, sc, var->Access()) + : ctx.dst->ty.pointer(param_type, sc); auto* param = ctx.dst->Param(new_var_symbol, param_type, attributes); ctx.InsertFront(func->params, param); is_pointer = true; @@ -228,7 +229,6 @@ struct ModuleScopeVarToEntryPointParam::State { const sem::Variable* var, Symbol new_var_symbol, bool& is_pointer) { - auto* var_ast = var->Declaration()->As(); auto* ty = var->Type()->UnwrapRef(); auto param_type = CreateASTTypeFor(ctx, ty); auto sc = var->AddressSpace(); @@ -254,7 +254,9 @@ struct ModuleScopeVarToEntryPointParam::State { // Use a pointer for non-handle types. utils::Vector attributes; if (!ty->is_handle()) { - param_type = ctx.dst->ty.pointer(param_type, sc, var_ast->declared_access); + param_type = sc == type::AddressSpace::kStorage + ? ctx.dst->ty.pointer(param_type, sc, var->Access()) + : ctx.dst->ty.pointer(param_type, sc); is_pointer = true; // Disable validation of the parameter's address space and of arguments passed to it. diff --git a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc index 999761dd49..0f4c5b9d2d 100644 --- a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc +++ b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc @@ -435,7 +435,7 @@ struct S { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { _ = *(tint_symbol); _ = *(tint_symbol_1); } @@ -465,7 +465,7 @@ struct S { auto* expect = R"( @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { _ = *(tint_symbol); _ = *(tint_symbol_1); } @@ -497,7 +497,7 @@ struct tint_symbol_1 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } )"; @@ -524,7 +524,7 @@ struct tint_symbol_1 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } )"; @@ -554,12 +554,12 @@ struct tint_symbol_2 { arr : array, } -fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr>) { +fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, read>) { _ = (*(tint_symbol))[0]; } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { foo(&((*(tint_symbol_1)).arr)); } )"; @@ -589,11 +589,11 @@ struct tint_symbol_2 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr) { foo(&((*(tint_symbol_1)).arr)); } -fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr>) { +fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, read>) { _ = (*(tint_symbol))[0]; } )"; @@ -624,7 +624,7 @@ struct tint_symbol_1 { alias myarray = array; @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } )"; @@ -652,7 +652,7 @@ struct tint_symbol_1 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } @@ -689,7 +689,7 @@ struct tint_symbol_1 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } )"; @@ -723,7 +723,7 @@ struct tint_symbol_1 { } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr) { _ = (*(tint_symbol)).arr[0]; } )"; @@ -773,12 +773,12 @@ struct S { fn no_uses() { } -fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { +fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { _ = *(tint_symbol); _ = *(tint_symbol_1); } -fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { +fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { let b : f32 = 2.0; _ = *(tint_symbol_2); bar(a, b, tint_symbol_2, tint_symbol_3); @@ -786,7 +786,7 @@ fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(di } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr) { foo(1.0, tint_symbol_4, tint_symbol_5); } )"; @@ -830,11 +830,11 @@ var s : S; auto* expect = R"( @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr) { +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr) { foo(1.0, tint_symbol_4, tint_symbol_5); } -fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { +fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { let b : f32 = 2.0; _ = *(tint_symbol_2); bar(a, b, tint_symbol_2, tint_symbol_3); @@ -844,7 +844,7 @@ fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(di fn no_uses() { } -fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { +fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { _ = *(tint_symbol); _ = *(tint_symbol_1); } diff --git a/src/tint/transform/renamer_test.cc b/src/tint/transform/renamer_test.cc index c233267a21..2d13436ccd 100644 --- a/src/tint/transform/renamer_test.cc +++ b/src/tint/transform/renamer_test.cc @@ -1927,7 +1927,9 @@ std::vector Identifiers() { out.push_back(ident); } for (auto* ident : type::kAddressSpaceStrings) { - out.push_back(ident); + if (!utils::HasPrefix(ident, "_")) { + out.push_back(ident); + } } for (auto* ident : type::kTexelFormatStrings) { out.push_back(ident); @@ -1940,24 +1942,24 @@ std::vector Identifiers() { using RenamerBuiltinIdentifierTest = TransformTestWithParam; -TEST_P(RenamerBuiltinIdentifierTest, GlobalVarName) { +TEST_P(RenamerBuiltinIdentifierTest, GlobalConstName) { auto expand = [&](const char* source) { return utils::ReplaceAll(source, "$name", GetParam()); }; auto src = expand(R"( -var $name = 42; +const $name = 42; fn f() { - var v = $name; + const v = $name; } )"); auto expect = expand(R"( -var tint_symbol = 42; +const tint_symbol = 42; fn tint_symbol_1() { - var tint_symbol_2 = tint_symbol; + const tint_symbol_2 = tint_symbol; } )"); diff --git a/src/tint/transform/std140.cc b/src/tint/transform/std140.cc index 4c487d0590..a19d78f73c 100644 --- a/src/tint/transform/std140.cc +++ b/src/tint/transform/std140.cc @@ -350,8 +350,8 @@ struct Std140::State { void ReplaceUniformVarTypes() { for (auto* global : src->AST().GlobalVariables()) { if (auto* var = global->As()) { - if (var->declared_address_space == type::AddressSpace::kUniform) { - auto* v = sem.Get(var); + auto* v = sem.Get(var); + if (v->AddressSpace() == type::AddressSpace::kUniform) { if (auto std140_ty = Std140Type(v->Type()->UnwrapRef())) { ctx.Replace(global->type.expr, b.Expr(std140_ty)); std140_uniforms.Add(v); diff --git a/src/tint/transform/zero_init_workgroup_memory.cc b/src/tint/transform/zero_init_workgroup_memory.cc index 49d96d901f..015a25bfd7 100644 --- a/src/tint/transform/zero_init_workgroup_memory.cc +++ b/src/tint/transform/zero_init_workgroup_memory.cc @@ -36,7 +36,8 @@ namespace { bool ShouldRun(const Program* program) { for (auto* global : program->AST().GlobalVariables()) { if (auto* var = global->As()) { - if (var->declared_address_space == type::AddressSpace::kWorkgroup) { + auto* v = program->Sem().Get(var); + if (v->AddressSpace() == type::AddressSpace::kWorkgroup) { return true; } } diff --git a/src/tint/type/address_space.cc b/src/tint/type/address_space.cc index 68f2c71f84..c3cfcc67ed 100644 --- a/src/tint/type/address_space.cc +++ b/src/tint/type/address_space.cc @@ -28,6 +28,12 @@ namespace tint::type { /// @param str the string to parse /// @returns the parsed enum, or AddressSpace::kUndefined if the string could not be parsed. AddressSpace ParseAddressSpace(std::string_view str) { + if (str == "__in") { + return AddressSpace::kIn; + } + if (str == "__out") { + return AddressSpace::kOut; + } if (str == "function") { return AddressSpace::kFunction; } @@ -53,14 +59,14 @@ std::ostream& operator<<(std::ostream& out, AddressSpace value) { switch (value) { case AddressSpace::kUndefined: return out << "undefined"; + case AddressSpace::kIn: + return out << "__in"; + case AddressSpace::kOut: + return out << "__out"; case AddressSpace::kFunction: return out << "function"; case AddressSpace::kHandle: return out << "handle"; - case AddressSpace::kIn: - return out << "in"; - case AddressSpace::kOut: - return out << "out"; case AddressSpace::kPrivate: return out << "private"; case AddressSpace::kPushConstant: diff --git a/src/tint/type/address_space.h b/src/tint/type/address_space.h index 5b3855c2c9..f0b0224591 100644 --- a/src/tint/type/address_space.h +++ b/src/tint/type/address_space.h @@ -30,10 +30,10 @@ namespace tint::type { /// Address space of a given pointer. enum class AddressSpace { kUndefined, + kIn, + kOut, kFunction, kHandle, // Tint-internal enum entry - not parsed - kIn, // Tint-internal enum entry - not parsed - kOut, // Tint-internal enum entry - not parsed kPrivate, kPushConstant, kStorage, @@ -52,7 +52,7 @@ std::ostream& operator<<(std::ostream& out, AddressSpace value); AddressSpace ParseAddressSpace(std::string_view str); constexpr const char* kAddressSpaceStrings[] = { - "function", "private", "push_constant", "storage", "uniform", "workgroup", + "__in", "__out", "function", "private", "push_constant", "storage", "uniform", "workgroup", }; /// @returns true if the AddressSpace is host-shareable diff --git a/src/tint/type/address_space_bench.cc b/src/tint/type/address_space_bench.cc index ecd3ccf57b..6f33497ee2 100644 --- a/src/tint/type/address_space_bench.cc +++ b/src/tint/type/address_space_bench.cc @@ -31,15 +31,62 @@ namespace { void AddressSpaceParser(::benchmark::State& state) { const char* kStrings[] = { - "fccnctin", "ucti3", "functVon", "function", "1unction", - "unJtqqon", "llun77tion", "ppqqivtHH", "prcv", "bivaGe", - "private", "priviive", "8WWivate", "pxxvate", "pXh_cggnstant", - "pX_Vonstanu", "push_consta3t", "push_constant", "push_constanE", "push_TTPnstant", - "puxxdh_constan", "s44orage", "stSSraVVe", "RtoR22e", "storage", - "sFra9e", "stoage", "VOORRHge", "unfoym", "llnnrrf77rm", - "unif4r00", "uniform", "nfoom", "zzform", "uiiippo1", - "workgrouXX", "wor55gro99nII", "wrrrkgroSSaHH", "workgroup", "kkrHoup", - "jgkrouRR", "wokroub", + "ccin", + "3", + "_Vin", + "__in", + "1_in", + "_qiJ", + "_lli77", + "__qHupp", + "vt", + "G_bt", + "__out", + "__viut", + "__8WWt", + "Mxxou", + "fuXggton", + "fuXtou", + "funct3on", + "function", + "funEtion", + "PPncTTion", + "xxuncddon", + "p44ivate", + "prSSvaVVe", + "RriR22e", + "private", + "pFva9e", + "priate", + "VOORRHte", + "push_constyn", + "punnh_crr77stallt", + "pu4h_cons00ant", + "push_constant", + "puoo_costan", + "ushzzcnstant", + "push_coii11apt", + "storaXXe", + "9II5tnnrage", + "stoaSSrHHYe", + "storage", + "stkke", + "jtogRa", + "sbrag", + "unifojm", + "niform", + "qform", + "uniform", + "uniNNrm", + "nifrvv", + "QQiform", + "workrorf", + "workjroup", + "wNNorkrou2", + "workgroup", + "workgrop", + "rrorkgroup", + "workgroGp", }; for (auto _ : state) { for (auto* str : kStrings) { diff --git a/src/tint/type/address_space_test.cc b/src/tint/type/address_space_test.cc index 8b2af5c661..7909f0ec45 100644 --- a/src/tint/type/address_space_test.cc +++ b/src/tint/type/address_space_test.cc @@ -42,6 +42,8 @@ inline std::ostream& operator<<(std::ostream& out, Case c) { } static constexpr Case kValidCases[] = { + {"__in", AddressSpace::kIn}, + {"__out", AddressSpace::kOut}, {"function", AddressSpace::kFunction}, {"private", AddressSpace::kPrivate}, {"push_constant", AddressSpace::kPushConstant}, @@ -51,15 +53,18 @@ static constexpr Case kValidCases[] = { }; static constexpr Case kInvalidCases[] = { - {"fccnctin", AddressSpace::kUndefined}, {"ucti3", AddressSpace::kUndefined}, - {"functVon", AddressSpace::kUndefined}, {"priv1te", AddressSpace::kUndefined}, - {"pqiJate", AddressSpace::kUndefined}, {"privat7ll", AddressSpace::kUndefined}, - {"pqqsh_pponstHnt", AddressSpace::kUndefined}, {"pus_cnstat", AddressSpace::kUndefined}, - {"bus_Gonstant", AddressSpace::kUndefined}, {"storiive", AddressSpace::kUndefined}, - {"8WWorage", AddressSpace::kUndefined}, {"sxxrage", AddressSpace::kUndefined}, - {"uXforgg", AddressSpace::kUndefined}, {"nfoXm", AddressSpace::kUndefined}, - {"unif3rm", AddressSpace::kUndefined}, {"workgroEp", AddressSpace::kUndefined}, - {"woTTPkroup", AddressSpace::kUndefined}, {"ddorkroxxp", AddressSpace::kUndefined}, + {"ccin", AddressSpace::kUndefined}, {"3", AddressSpace::kUndefined}, + {"_Vin", AddressSpace::kUndefined}, {"__ou1", AddressSpace::kUndefined}, + {"qq_Jt", AddressSpace::kUndefined}, {"__oll7t", AddressSpace::kUndefined}, + {"qquntppHon", AddressSpace::kUndefined}, {"cnciv", AddressSpace::kUndefined}, + {"funGion", AddressSpace::kUndefined}, {"priviive", AddressSpace::kUndefined}, + {"8WWivate", AddressSpace::kUndefined}, {"pxxvate", AddressSpace::kUndefined}, + {"pXh_cggnstant", AddressSpace::kUndefined}, {"pX_Vonstanu", AddressSpace::kUndefined}, + {"push_consta3t", AddressSpace::kUndefined}, {"Etorage", AddressSpace::kUndefined}, + {"sPTTrage", AddressSpace::kUndefined}, {"storadxx", AddressSpace::kUndefined}, + {"u44iform", AddressSpace::kUndefined}, {"unSSfoVVm", AddressSpace::kUndefined}, + {"RniR22m", AddressSpace::kUndefined}, {"w9rFroup", AddressSpace::kUndefined}, + {"workgoup", AddressSpace::kUndefined}, {"woVROOrHup", AddressSpace::kUndefined}, }; using AddressSpaceParseTest = testing::TestWithParam; diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc index ae34abb7dc..c92e51ad6d 100644 --- a/src/tint/writer/wgsl/generator_impl.cc +++ b/src/tint/writer/wgsl/generator_impl.cc @@ -34,8 +34,6 @@ #include "src/tint/ast/workgroup_attribute.h" #include "src/tint/sem/struct.h" #include "src/tint/sem/switch_statement.h" -#include "src/tint/type/access.h" -#include "src/tint/type/texture_dimension.h" #include "src/tint/utils/math.h" #include "src/tint/utils/scoped_assignment.h" #include "src/tint/writer/float_to_string.h" @@ -376,24 +374,6 @@ bool GeneratorImpl::EmitImageFormat(std::ostream& out, const type::TexelFormat f return true; } -bool GeneratorImpl::EmitAccess(std::ostream& out, const type::Access access) { - switch (access) { - case type::Access::kRead: - out << "read"; - return true; - case type::Access::kWrite: - out << "write"; - return true; - case type::Access::kReadWrite: - out << "read_write"; - return true; - default: - break; - } - diagnostics_.add_error(diag::System::Writer, "unknown access"); - return false; -} - bool GeneratorImpl::EmitStructType(const ast::Struct* str) { if (str->attributes.Length()) { if (!EmitAttributes(line(), str->attributes)) { @@ -473,17 +453,18 @@ bool GeneratorImpl::EmitVariable(std::ostream& out, const ast::Variable* v) { v, // [&](const ast::Var* var) { out << "var"; - auto address_space = var->declared_address_space; - auto ac = var->declared_access; - if (address_space != type::AddressSpace::kUndefined || ac != type::Access::kUndefined) { - out << "<" << address_space; - if (ac != type::Access::kUndefined) { + if (var->declared_address_space || var->declared_access) { + out << "<"; + TINT_DEFER(out << ">"); + if (!EmitExpression(out, var->declared_address_space)) { + return false; + } + if (var->declared_access) { out << ", "; - if (!EmitAccess(out, ac)) { + if (!EmitExpression(out, var->declared_access)) { return false; } } - out << ">"; } return true; }, diff --git a/src/tint/writer/wgsl/generator_impl.h b/src/tint/writer/wgsl/generator_impl.h index 56c0dfb75d..746ce0f798 100644 --- a/src/tint/writer/wgsl/generator_impl.h +++ b/src/tint/writer/wgsl/generator_impl.h @@ -35,7 +35,6 @@ #include "src/tint/ast/unary_op_expression.h" #include "src/tint/program.h" #include "src/tint/sem/struct.h" -#include "src/tint/type/storage_texture.h" #include "src/tint/writer/text_generator.h" namespace tint::writer::wgsl { @@ -209,11 +208,6 @@ class GeneratorImpl : public TextGenerator { /// @param fmt the format to generate /// @returns true if the format is emitted bool EmitImageFormat(std::ostream& out, const type::TexelFormat fmt); - /// Handles emitting an access control - /// @param out the output stream - /// @param access the access to generate - /// @returns true if the access is emitted - bool EmitAccess(std::ostream& out, const type::Access access); /// Handles a unary op expression /// @param out the output stream /// @param expr the expression to emit