tint: Unkeyword 'var' template args
Change the address space and access mode in ast::Var from enums
to Expressions. Have the resolver resolve these, like we do for
other template arguments.
As the AST nodes now have identifier expressions, the tint-internal
'in' and 'out' address spaces have been prefixed with underscores to
prevent input code from using this.
Change-Id: Ie8abf371ee6a7031613709b83b575d2723418fcf
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/120405
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index dd679c7..fbed83a 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1646,7 +1646,6 @@
"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 67ca5e9..4b87d7b 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -1092,7 +1092,6 @@
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 b748381..bccdbcc 100644
--- a/src/tint/ast/var.cc
+++ b/src/tint/ast/var.cc
@@ -25,8 +25,8 @@
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<const Attribute*> attrs)
: Base(pid, nid, src, n, ty, init, std::move(attrs)),
@@ -45,10 +45,11 @@
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<Var>(src, n, ty, declared_address_space, declared_access, init,
- std::move(attrs));
+ return ctx->dst->create<Var>(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 e3537d1..006d8de 100644
--- a/src/tint/ast/var.h
+++ b/src/tint/ast/var.h
@@ -56,8 +56,8 @@
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<const Attribute*> attributes);
@@ -77,10 +77,10 @@
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 88bcdc4..dcbf1e5 100644
--- a/src/tint/ast/variable_test.cc
+++ b/src/tint/ast/variable_test.cc
@@ -28,7 +28,8 @@
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 @@
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 @@
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 ee28010..f55b12c 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -80,9 +80,9 @@
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 1c04faf..f89a30b 100644
--- a/src/tint/program_builder.h
+++ b/src/tint/program_builder.h
@@ -219,24 +219,34 @@
/// constructing an ast::Var.
struct VarOptions {
template <typename... ARGS>
- explicit VarOptions(ARGS&&... args) {
- (Set(std::forward<ARGS>(args)), ...);
+ explicit VarOptions(ProgramBuilder& b, ARGS&&... args) {
+ (Set(b, std::forward<ARGS>(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<const ast::Attribute*, 4> 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<const ast::Attribute*> 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<const ast::Attribute*> 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 @@
/// @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 @@
/// @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 <typename NAME, typename... OPTIONS>
const ast::Var* Var(const Source& source, NAME&& name, OPTIONS&&... options) {
- VarOptions opts(std::forward<OPTIONS>(options)...);
+ VarOptions opts(*this, std::forward<OPTIONS>(options)...);
return create<ast::Var>(source, Ident(std::forward<NAME>(name)), opts.type,
opts.address_space, opts.access, opts.initializer,
std::move(opts.attributes));
@@ -2091,8 +2101,7 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 @@
/// @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 046780f..513b7f3 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -2522,7 +2522,8 @@
}
}
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<ast::VariableDeclStatement>(Source{}, var);
AddStatement(var_decl_stmt);
auto* var_type = ty_.Reference(var_store_type, type::AddressSpace::kUndefined);
@@ -3367,8 +3368,9 @@
// no need to remap pointer properties.
auto* store_type = parser_impl_.ConvertType(def_inst->type_id());
AddStatement(create<ast::VariableDeclStatement>(
- 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 @@
// 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 5707bf1..443b28f 100644
--- a/src/tint/reader/spirv/parser_impl.cc
+++ b/src/tint/reader/spirv/parser_impl.cc
@@ -1504,12 +1504,15 @@
// 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 @@
<< 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 @@
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<Named>()) {
+ 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 @@
return nullptr;
}
- type::Access access = type::Access::kUndefined;
- if (address_space == type::AddressSpace::kStorage) {
- bool read_only = false;
- if (auto* tn = storage_type->As<Named>()) {
- 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 5d6daa4..0c77a02 100644
--- a/src/tint/reader/spirv/parser_impl.h
+++ b/src/tint/reader/spirv/parser_impl.h
@@ -420,10 +420,17 @@
/// @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 @@
/// 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 @@
/// 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 @@
std::unordered_map<const spvtools::opt::Instruction*, const Type*> handle_type_;
/// Maps the SPIR-V ID of a module-scope variable to its AST variable.
- utils::Hashmap<uint32_t, const ast::Var*, 16> module_variable_;
+ utils::Hashmap<uint32_t, ModuleVariable, 16> 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 0ccef95..79d39d2 100644
--- a/src/tint/reader/wgsl/classify_template_args.cc
+++ b/src/tint/reader/wgsl/classify_template_args.cc
@@ -70,8 +70,8 @@
for (size_t i = 0; i < count - 1; i++) {
switch (tokens[i].type()) {
- // <identifier> + 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 c503353..c9a8576 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -204,23 +204,6 @@
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 @@
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<ast::Var>(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 @@
return expect_ident_with_optional_type_specifier(use, /* allow_inferred */ false);
}
-// access_mode
-// : 'read'
-// | 'write'
-// | 'read_write'
-Expect<type::Access> 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::VariableQualifier> 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<VariableQualifier> {
- auto source = make_source_range();
- auto sc = expect_address_space(use);
- if (sc.errored) {
+ auto vq = expect_template_arg_block(use, [&]() -> Expect<VariableQualifier> {
+ 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>{VariableQualifier{sc.value, type::Access::kUndefined},
- source};
+ return VariableQualifier{address_space.value};
});
if (vq.errored) {
@@ -900,18 +874,6 @@
return type.value;
}
-// address_space
-// : 'function'
-// | 'private'
-// | 'workgroup'
-// | 'uniform'
-// | 'storage'
-//
-// Note, we also parse `push_constant` from the experimental extension
-Expect<type::AddressSpace> 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<const ast::Struct*> ParserImpl::struct_decl() {
@@ -1519,12 +1481,13 @@
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<ast::Var>(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<ast::VariableDeclStatement>(var->source, var);
}
@@ -2520,7 +2483,7 @@
return create<ast::BinaryExpression>(tok_op.source(), op, lhs, rhs.value);
}
-Expect<const ast::Expression*> ParserImpl::expect_expression() {
+Expect<const ast::Expression*> ParserImpl::expect_expression(std::string_view use) {
auto& t = peek();
auto expr = expression();
if (expr.errored) {
@@ -2529,7 +2492,7 @@
if (expr.matched) {
return expr.value;
}
- return add_error(t, "expected expression");
+ return add_error(t, "expected expression for " + std::string(use));
}
Expect<utils::Vector<const ast::Expression*, 3>> ParserImpl::expect_expression_list(
@@ -2537,7 +2500,7 @@
Token::Type terminator) {
utils::Vector<const ast::Expression*, 3> 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 f692d77..64618b3 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -148,10 +148,9 @@
/// Constructor for a successful parse.
/// @param val the result value of the parse
- /// @param s the optional source of the value
template <typename U>
- inline Maybe(U&& val, const Source& s = {}) // NOLINT
- : value(std::forward<U>(val)), source(s), matched(true) {}
+ inline Maybe(U&& val) // NOLINT
+ : value(std::forward<U>(val)), matched(true) {}
/// Constructor for parse error state.
inline Maybe(Failure::Errored) : errored(true) {} // NOLINT
@@ -163,16 +162,13 @@
/// @param e the Expect to copy this Maybe from
template <typename U>
inline Maybe(const Expect<U>& 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 <typename U>
inline Maybe(Expect<U>&& 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 @@
/// 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 @@
/// 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 @@
/// 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 @@
/// Parses a `type_specifier` grammar element
/// @returns the parsed Type or nullptr if none matched.
Maybe<ast::Type> 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<type::AddressSpace> expect_address_space(std::string_view use);
/// Parses a `struct_decl` grammar element.
/// @returns the struct type or nullptr on error
Maybe<const ast::Struct*> struct_decl();
@@ -482,11 +453,6 @@
/// not match a stage name.
/// @returns the pipeline stage.
Expect<ast::PipelineStage> 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<type::Access> 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 @@
/// @returns the parsed expression or nullptr
Maybe<const ast::Expression*> expression();
/// Parses the `expression` grammar rule
+ /// @param use the use of the expression
/// @returns the parsed expression or error
- Expect<const ast::Expression*> expect_expression();
+ Expect<const ast::Expression*> 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 32c3c47..0000000
--- 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<AddressSpaceData> {};
-
-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 d40898a..18b40be 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 @@
TEST_F(ParserImplErrorTest, GlobalDeclStorageTextureMissingSubtype) {
EXPECT("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 @@
)");
}
-TEST_F(ParserImplErrorTest, GlobalDeclVarStorageDeclInvalidClass) {
- EXPECT("var<fish> i : i32",
- R"(test.wgsl:1:5 error: expected address space for variable declaration
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup'
-var<fish> i : i32
- ^^^^
-)");
-}
-
TEST_F(ParserImplErrorTest, GlobalDeclVarStorageDeclMissingGThan) {
EXPECT("var<private i : i32",
- R"(test.wgsl:1:13 error: expected '>' for variable declaration
+ R"(test.wgsl:1:4 error: missing closing '>' for variable declaration
var<private i : i32
- ^
+ ^
)");
}
diff --git a/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
index 92272cc..daf978e 100644
--- a/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
@@ -31,10 +31,8 @@
ASSERT_NE(var, nullptr);
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);
@@ -58,8 +56,7 @@
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 @@
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 @@
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 @@
EXPECT_EQ(p->error(), "1:24: missing initializer for 'var' declaration");
}
-TEST_F(ParserImplTest, GlobalVariableDecl_InvalidVariableDecl) {
- auto p = parser("var<invalid> 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 b40d887..25e60e0 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 @@
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 @@
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 @@
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 @@
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 @@
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 9138d55..993ecb0 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 @@
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 @@
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<unknown> 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 111f207..9a8b936 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 @@
TEST_P(VariableQualifierTest, ParsesAddressSpace) {
auto params = GetParam();
- auto p = parser(std::string("<") + params.input + ">");
+ 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 @@
VariableStorageData{"storage, read_write", type::AddressSpace::kStorage,
type::Access::kReadWrite}));
-TEST_F(ParserImplTest, VariableQualifier_NoMatch) {
- auto p = parser("<not-a-storage-class>");
- 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 @@
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 2c879de..95301c7 100644
--- a/src/tint/resolver/address_space_validation_test.cc
+++ b/src/tint/resolver/address_space_validation_test.cc
@@ -32,8 +32,9 @@
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 @@
EXPECT_EQ(
r()->error(),
- R"(12:34 error: only variables in <storage> address space may declare an access mode)");
+ R"(12:34 error: only variables in <storage> address space may specify an access mode)");
}
TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NotStorage_AccessMode) {
@@ -481,7 +482,7 @@
EXPECT_EQ(
r()->error(),
- R"(12:34 error: only pointers in <storage> address space may declare an access mode)");
+ R"(12:34 error: only pointers in <storage> 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 16a3f8e..75e377f 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -183,13 +183,9 @@
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 @@
}
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<ast::Var>()) {
+ 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 @@
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 @@
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 31e484e..a316ad0 100644
--- a/src/tint/resolver/dependency_graph_test.cc
+++ b/src/tint/resolver/dependency_graph_test.cc
@@ -1437,13 +1437,13 @@
<< 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 b04831e..1062958 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -1088,15 +1088,14 @@
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 @@
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 b87852f..8336765 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -473,8 +473,14 @@
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 @@
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 73fa9b8..8ee795c 100644
--- a/src/tint/resolver/validation_test.cc
+++ b/src/tint/resolver/validation_test.cc
@@ -295,22 +295,23 @@
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<f32>' must not have a address space)");
+ EXPECT_EQ(
+ r()->error(),
+ R"(12:34 error: variables of type 'texture_1d<f32>' 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 5b399cf..2935c70 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -299,7 +299,7 @@
// * 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 <storage> address space may declare an access mode",
+ AddError("only pointers in <storage> address space may specify an access mode",
a->source);
return false;
}
@@ -604,7 +604,7 @@
}
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 @@
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 @@
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 <storage> address space may declare an access mode",
+ if (v->AddressSpace() != type::AddressSpace::kStorage) {
+ AddError("only variables in <storage> address space may specify an access mode",
var->source);
return false;
}
@@ -726,8 +727,8 @@
}
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 b657980..192b2c9 100644
--- a/src/tint/transform/binding_remapper.cc
+++ b/src/tint/transform/binding_remapper.cc
@@ -122,11 +122,12 @@
// 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<uint32_t>(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<uint32_t>(access)) +
+ ")");
return Program(std::move(b));
}
auto* sem = src->Sem().Get(var);
@@ -139,9 +140,14 @@
}
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<ast::Var>(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 5af8a14..f0ce8c3 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 @@
)";
auto* expect = R"(
-@location(1) @internal(disable_validation__ignore_address_space) var<in> 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<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
var col : f32 = (coord.x * loc1);
@@ -251,13 +251,13 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
struct FragBuiltins {
coord : vec4<f32>,
@@ -304,13 +304,13 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
fn frag_main_inner(loc0 : f32, locations : FragLocations, builtins : FragBuiltins) {
var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
@@ -567,7 +567,7 @@
)";
auto* expect = R"(
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> 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 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<out> color_1 : vec4<f32>;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4<f32>;
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> 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<out> mask_1 : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array<u32, 1u>;
struct FragOutput {
color : vec4<f32>,
@@ -729,11 +729,11 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<out> color_1 : vec4<f32>;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4<f32>;
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> 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<out> mask_1 : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array<u32, 1u>;
fn frag_main_inner() -> FragOutput {
var output : FragOutput;
@@ -1028,13 +1028,13 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> 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<in> mul_2 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_2 : f32;
struct FragmentInput {
value : f32,
@@ -1094,13 +1094,13 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> 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<in> 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 @@
auto* expect =
R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> vi_1 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4<i32>;
-@location(3) @internal(disable_validation__ignore_address_space) var<in> vu_1 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4<u32>;
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> 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<out> 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<out> vi_2 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4<i32>;
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vu_2 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4<u32>;
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> vi_3 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4<i32>;
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vu_3 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4<u32>;
-@location(0) @internal(disable_validation__ignore_address_space) var<out> 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<out> 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<out> vi_4 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4<i32>;
-@location(3) @internal(disable_validation__ignore_address_space) var<out> vu_4 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4<u32>;
struct VertexIn {
i : i32,
@@ -2082,39 +2082,39 @@
auto* expect =
R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> vi_1 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4<i32>;
-@location(3) @internal(disable_validation__ignore_address_space) var<in> vu_1 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4<u32>;
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> 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<out> 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<out> vi_2 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4<i32>;
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vu_2 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4<u32>;
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<in> vi_3 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4<i32>;
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vu_3 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4<u32>;
-@location(0) @internal(disable_validation__ignore_address_space) var<out> 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<out> 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<out> vi_4 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4<i32>;
-@location(3) @internal(disable_validation__ignore_address_space) var<out> vu_4 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4<u32>;
fn vert_main_inner(in : VertexIn) -> VertexOut {
return VertexOut(in.i, in.u, in.vi, in.vu, vec4<f32>());
@@ -3161,9 +3161,9 @@
)";
auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> value : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> value : vec4<f32>;
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32;
fn vert_main_inner() -> vec4<f32> {
return vec4<f32>();
@@ -3236,9 +3236,9 @@
)";
auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32;
struct VertOut {
pos : vec4<f32>,
@@ -3277,9 +3277,9 @@
)";
auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> 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 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<out> 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<out> vertex_point_size_1_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4<f32>;
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_4 : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_4 : f32;
var<private> vertex_point_size : f32;
@@ -3502,15 +3502,15 @@
)";
auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<out> 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<out> vertex_point_size_1_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4<f32>;
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> 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 @@
)";
auto* expect = R"(
-@builtin(sample_index) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> 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<in> mask_in_1 : array<u32, 1u>;
+@builtin(sample_mask) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> mask_in_1 : array<u32, 1u>;
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> value : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> value : array<u32, 1u>;
fn main_inner(sample_index : u32, mask_in : u32) -> u32 {
return mask_in;
@@ -3903,11 +3903,11 @@
)";
auto* expect = R"(
-@builtin(sample_index) @internal(disable_validation__ignore_address_space) var<in> 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<in> gl_SampleMaskIn : array<i32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__in> gl_SampleMaskIn : array<i32, 1u>;
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> gl_SampleMask : array<i32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> gl_SampleMask : array<i32, 1u>;
fn fragment_main(sample_index : u32, mask_in : u32) -> u32 {
return mask_in;
@@ -3938,11 +3938,11 @@
)";
auto* expect = R"(
-@builtin(vertex_index) @internal(disable_validation__ignore_address_space) var<in> 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<in> 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<out> gl_Position : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> gl_Position : vec4<f32>;
fn vertex_main(vertexID : u32, instanceID : u32) -> vec4<f32> {
return vec4<f32>((f32(vertexID) + f32(instanceID)));
diff --git a/src/tint/transform/clamp_frag_depth.cc b/src/tint/transform/clamp_frag_depth.cc
index 57d9a40..8d84ed3 100644
--- a/src/tint/transform/clamp_frag_depth.cc
+++ b/src/tint/transform/clamp_frag_depth.cc
@@ -88,10 +88,11 @@
// Abort on any use of push constants in the module.
for (auto* global : src->AST().GlobalVariables()) {
if (auto* var = global->As<ast::Var>()) {
- 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 ff65fec..47356e9 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 @@
WorkgroupParameterMemberList& workgroup_parameter_members,
bool& is_pointer,
bool& is_wrapped) {
- auto* var_ast = var->Declaration()->As<ast::Var>();
auto* ty = var->Type()->UnwrapRef();
// Helper to create an AST node for the store type of the variable.
@@ -160,7 +159,9 @@
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 @@
const sem::Variable* var,
Symbol new_var_symbol,
bool& is_pointer) {
- auto* var_ast = var->Declaration()->As<ast::Var>();
auto* ty = var->Type()->UnwrapRef();
auto param_type = CreateASTTypeFor(ctx, ty);
auto sc = var->AddressSpace();
@@ -254,7 +254,9 @@
// Use a pointer for non-handle types.
utils::Vector<const ast::Attribute*, 2> 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 999761d..0f4c5b9 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 @@
}
@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<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S, read>) {
_ = *(tint_symbol);
_ = *(tint_symbol_1);
}
@@ -465,7 +465,7 @@
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<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S, read>) {
_ = *(tint_symbol);
_ = *(tint_symbol_1);
}
@@ -497,7 +497,7 @@
}
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
)";
@@ -524,7 +524,7 @@
}
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
)";
@@ -554,12 +554,12 @@
arr : array<f32>,
}
-fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>>) {
+fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>, 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<storage, tint_symbol_2>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2, read>) {
foo(&((*(tint_symbol_1)).arr));
}
)";
@@ -589,11 +589,11 @@
}
@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<storage, tint_symbol_2>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2, read>) {
foo(&((*(tint_symbol_1)).arr));
}
-fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>>) {
+fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>, read>) {
_ = (*(tint_symbol))[0];
}
)";
@@ -624,7 +624,7 @@
alias myarray = array<f32>;
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
)";
@@ -652,7 +652,7 @@
}
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -689,7 +689,7 @@
}
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
)";
@@ -723,7 +723,7 @@
}
@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<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
_ = (*(tint_symbol)).arr[0];
}
)";
@@ -773,12 +773,12 @@
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<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S>) {
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S, read>) {
_ = *(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<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S>) {
+fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S, read>) {
let b : f32 = 2.0;
_ = *(tint_symbol_2);
bar(a, b, tint_symbol_2, tint_symbol_3);
@@ -786,7 +786,7 @@
}
@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<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S, read>) {
foo(1.0, tint_symbol_4, tint_symbol_5);
}
)";
@@ -830,11 +830,11 @@
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<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S, read>) {
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<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S>) {
+fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S, read>) {
let b : f32 = 2.0;
_ = *(tint_symbol_2);
bar(a, b, tint_symbol_2, tint_symbol_3);
@@ -844,7 +844,7 @@
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<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S>) {
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S, read>) {
_ = *(tint_symbol);
_ = *(tint_symbol_1);
}
diff --git a/src/tint/transform/renamer_test.cc b/src/tint/transform/renamer_test.cc
index c233267..2d13436 100644
--- a/src/tint/transform/renamer_test.cc
+++ b/src/tint/transform/renamer_test.cc
@@ -1927,7 +1927,9 @@
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 @@
using RenamerBuiltinIdentifierTest = TransformTestWithParam<const char*>;
-TEST_P(RenamerBuiltinIdentifierTest, GlobalVarName) {
+TEST_P(RenamerBuiltinIdentifierTest, GlobalConstName) {
auto expand = [&](const char* source) {
return utils::ReplaceAll(source, "$name", GetParam());
};
auto src = expand(R"(
-var<private> $name = 42;
+const $name = 42;
fn f() {
- var v = $name;
+ const v = $name;
}
)");
auto expect = expand(R"(
-var<private> 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 4c487d0..a19d78f 100644
--- a/src/tint/transform/std140.cc
+++ b/src/tint/transform/std140.cc
@@ -350,8 +350,8 @@
void ReplaceUniformVarTypes() {
for (auto* global : src->AST().GlobalVariables()) {
if (auto* var = global->As<ast::Var>()) {
- 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 49d96d9..015a25b 100644
--- a/src/tint/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/transform/zero_init_workgroup_memory.cc
@@ -36,7 +36,8 @@
bool ShouldRun(const Program* program) {
for (auto* global : program->AST().GlobalVariables()) {
if (auto* var = global->As<ast::Var>()) {
- 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 68f2c71..c3cfcc6 100644
--- a/src/tint/type/address_space.cc
+++ b/src/tint/type/address_space.cc
@@ -28,6 +28,12 @@
/// @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 @@
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 5b3855c..f0b0224 100644
--- a/src/tint/type/address_space.h
+++ b/src/tint/type/address_space.h
@@ -30,10 +30,10 @@
/// 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 @@
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 ecd3ccf..6f33497 100644
--- a/src/tint/type/address_space_bench.cc
+++ b/src/tint/type/address_space_bench.cc
@@ -31,15 +31,62 @@
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 8b2af5c..7909f0e 100644
--- a/src/tint/type/address_space_test.cc
+++ b/src/tint/type/address_space_test.cc
@@ -42,6 +42,8 @@
}
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 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<Case>;
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index ae34abb..c92e51a 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 @@
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 @@
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 56c0dfb..746ce0f 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 @@
/// @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
