encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Ben Clayton 2023-02-18 17:13:18 +00:00 committed by Dawn LUCI CQ
parent 28ec26b678
commit 79781f26d1
40 changed files with 492 additions and 557 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/tint/BUILD.gn
View File

@ -1646,7 +1646,6 @@ if (tint_build_unittests) {
"reader/wgsl/classify_template_args_test.cc",
"reader/wgsl/lexer_test.cc",
"reader/wgsl/parser_impl_additive_expression_test.cc",
"reader/wgsl/parser_impl_address_space_test.cc",
"reader/wgsl/parser_impl_argument_expression_list_test.cc",
"reader/wgsl/parser_impl_assignment_stmt_test.cc",
"reader/wgsl/parser_impl_bitwise_expression_test.cc",

1
src/tint/CMakeLists.txt
View File

@ -1092,7 +1092,6 @@ if(TINT_BUILD_TESTS)
reader/wgsl/parser_impl_singular_expression_test.cc
reader/wgsl/parser_impl_statement_test.cc
reader/wgsl/parser_impl_statements_test.cc
reader/wgsl/parser_impl_address_space_test.cc
reader/wgsl/parser_impl_struct_body_decl_test.cc
reader/wgsl/parser_impl_struct_decl_test.cc
reader/wgsl/parser_impl_struct_attribute_decl_test.cc

9
src/tint/ast/var.cc
View File

@ -25,8 +25,8 @@ Var::Var(ProgramID pid,
const Source& src,
const Identifier* n,
Type ty,
type::AddressSpace address_space,
type::Access access,
const Expression* address_space,
const Expression* access,
const Expression* init,
utils::VectorRef<const Attribute*> attrs)
: Base(pid, nid, src, n, ty, init, std::move(attrs)),
@ -45,10 +45,11 @@ const Var* Var::Clone(CloneContext* ctx) const {
auto src = ctx->Clone(source);
auto* n = ctx->Clone(name);
auto ty = ctx->Clone(type);
auto* address_space = ctx->Clone(declared_address_space);
auto* access = ctx->Clone(declared_access);
auto* init = ctx->Clone(initializer);
auto attrs = ctx->Clone(attributes);
return ctx->dst->create<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

8
src/tint/ast/var.h
View File

@ -56,8 +56,8 @@ class Var final : public Castable<Var, Variable> {
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 @@ class Var final : public Castable<Var, Variable> {
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

7
src/tint/ast/variable_test.cc
View File

@ -28,7 +28,8 @@ TEST_F(VariableTest, Creation) {
auto* v = Var("my_var", ty.i32(), type::AddressSpace::kFunction);
CheckIdentifier(Symbols(), v->name, "my_var");
EXPECT_EQ(v->declared_address_space, type::AddressSpace::kFunction);
CheckIdentifier(Symbols(), v->declared_address_space, "function");
EXPECT_EQ(v->declared_access, nullptr);
CheckIdentifier(Symbols(), v->type, "i32");
EXPECT_EQ(v->source.range.begin.line, 0u);
EXPECT_EQ(v->source.range.begin.column, 0u);
@ -41,7 +42,7 @@ TEST_F(VariableTest, CreationWithSource) {
ty.f32(), type::AddressSpace::kPrivate, utils::Empty);
CheckIdentifier(Symbols(), v->name, "i");
EXPECT_EQ(v->declared_address_space, type::AddressSpace::kPrivate);
CheckIdentifier(Symbols(), v->declared_address_space, "private");
CheckIdentifier(Symbols(), v->type, "f32");
EXPECT_EQ(v->source.range.begin.line, 27u);
EXPECT_EQ(v->source.range.begin.column, 4u);
@ -54,7 +55,7 @@ TEST_F(VariableTest, CreationEmpty) {
ty.i32(), type::AddressSpace::kWorkgroup, utils::Empty);
CheckIdentifier(Symbols(), v->name, "a_var");
EXPECT_EQ(v->declared_address_space, type::AddressSpace::kWorkgroup);
CheckIdentifier(Symbols(), v->declared_address_space, "workgroup");
CheckIdentifier(Symbols(), v->type, "i32");
EXPECT_EQ(v->source.range.begin.line, 27u);
EXPECT_EQ(v->source.range.begin.column, 4u);

4
src/tint/intrinsics.def
View File

@ -80,9 +80,9 @@ enum address_space {
uniform
storage
push_constant
__in
__out
@internal handle
@internal in
@internal out
}
// https://gpuweb.github.io/gpuweb/wgsl/#memory-access-mode

63
src/tint/program_builder.h
View File

@ -219,24 +219,34 @@ class ProgramBuilder {
/// 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,8 +2065,8 @@ class ProgramBuilder {
/// @param name the variable name
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Type* - specifies the variable type
/// * type::AddressSpace - specifies the variable address space
/// * 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)
@ -2072,8 +2082,8 @@ class ProgramBuilder {
/// @param name the variable name
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Type* - specifies the variable type
/// * type::AddressSpace - specifies the variable address space
/// * 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)
@ -2081,7 +2091,7 @@ class ProgramBuilder {
/// @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 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Identifier* - specifies the variable type
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Const` with the given name, type and additional options
@ -2106,8 +2115,8 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Identifier* - specifies the variable type
/// * ast::Type* - specifies the variable type
/// * ast::Identifier* - specifies the variable's type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Const` with the given name, type and additional options
@ -2122,7 +2131,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Let` with the given name, type and additional options
@ -2136,7 +2145,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Let` with the given name, type and additional options
@ -2175,7 +2184,7 @@ class ProgramBuilder {
/// @param name the variable name
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * type::AddressSpace - specifies the variable address space
/// * type::Access - specifies the variable's access control
/// * ast::Expression* - specifies the variable's initializer expression
@ -2192,7 +2201,7 @@ class ProgramBuilder {
/// @param name the variable name
/// @param options the extra options passed to the ast::Var initializer
/// Can be any of the following, in any order:
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * type::AddressSpace - specifies the variable address space
/// * type::Access - specifies the variable's access control
/// * ast::Expression* - specifies the variable's initializer expression
@ -2211,7 +2220,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Const initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Const` with the given name, type and additional options, which is
@ -2226,7 +2235,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Const initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Const` with the given name, type and additional options, which is
@ -2242,7 +2251,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Override initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Override` with the given name, type and additional options, which is
@ -2257,7 +2266,7 @@ class ProgramBuilder {
/// @param options the extra options passed to the ast::Override initializer
/// Can be any of the following, in any order:
/// * ast::Expression* - specifies the variable's initializer expression (required)
/// * ast::Type* - specifies the variable type
/// * ast::Type - specifies the variable's type
/// * ast::Attribute* - specifies the variable's attributes (repeatable, or vector)
/// Note that non-repeatable arguments of the same type will use the last argument's value.
/// @returns an `ast::Override` with the given name, type and additional options, which is

13
src/tint/reader/spirv/function.cc
View File

@ -2522,7 +2522,8 @@ bool FunctionEmitter::EmitFunctionVariables() {
}
}
auto* var = parser_impl_.MakeVar(inst.result_id(), type::AddressSpace::kUndefined,
var_store_type, initializer, AttributeList{});
type::Access::kUndefined, var_store_type, initializer,
AttributeList{});
auto* var_decl_stmt = create<ast::VariableDeclStatement>(Source{}, var);
AddStatement(var_decl_stmt);
auto* var_type = ty_.Reference(var_store_type, type::AddressSpace::kUndefined);
@ -3367,8 +3368,9 @@ bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info,
// no need to remap pointer properties.
auto* store_type = parser_impl_.ConvertType(def_inst->type_id());
AddStatement(create<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 @@ DefInfo::Pointer FunctionEmitter::GetPointerInfo(uint32_t id) {
// either variables or function parameters.
switch (opcode(inst)) {
case spv::Op::OpVariable: {
if (const auto* module_var = parser_impl_.GetModuleVariable(id)) {
return DefInfo::Pointer{module_var->declared_address_space,
module_var->declared_access};
if (auto v = parser_impl_.GetModuleVariable(id); v.var) {
return DefInfo::Pointer{v.address_space, v.access};
}
// Local variables are always Function storage class, with default
// access mode.

49
src/tint/reader/spirv/parser_impl.cc
View File

@ -1504,12 +1504,15 @@ bool ParserImpl::EmitModuleScopeVariables() {
// here.)
ast_initializer = MakeConstantExpression(var.GetSingleWordInOperand(1)).expr;
}
auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_store_type, ast_initializer,
utils::Empty);
auto ast_access = VarAccess(ast_store_type, ast_address_space);
auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_access, ast_store_type,
ast_initializer, utils::Empty);
// TODO(dneto): initializers (a.k.a. initializer expression)
if (ast_var) {
builder_.AST().AddGlobalVariable(ast_var);
module_variable_.GetOrCreate(var.result_id(), [ast_var] { return ast_var; });
module_variable_.GetOrCreate(var.result_id(), [&] {
return ModuleVariable{ast_var, ast_address_space, ast_access};
});
}
}
@ -1536,14 +1539,16 @@ bool ParserImpl::EmitModuleScopeVariables() {
<< init->PrettyPrint();
}
}
auto* ast_var =
MakeVar(builtin_position_.per_vertex_var_id,
enum_converter_.ToAddressSpace(builtin_position_.storage_class),
ConvertType(builtin_position_.position_member_type_id), ast_initializer, {});
auto storage_type = ConvertType(builtin_position_.position_member_type_id);
auto ast_address_space = enum_converter_.ToAddressSpace(builtin_position_.storage_class);
auto ast_access = VarAccess(storage_type, ast_address_space);
auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space, ast_access,
storage_type, ast_initializer, {});
builder_.AST().AddGlobalVariable(ast_var);
module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id,
[ast_var] { return ast_var; });
module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id, [&] {
return ModuleVariable{ast_var, ast_address_space};
});
}
return success_;
}
@ -1571,25 +1576,29 @@ const spvtools::opt::analysis::IntConstant* ParserImpl::GetArraySize(uint32_t va
return size->AsIntConstant();
}
const ast::Var* ParserImpl::MakeVar(uint32_t id,
type::AddressSpace address_space,
const Type* storage_type,
const ast::Expression* initializer,
AttributeList decorations) {
if (storage_type == nullptr) {
Fail() << "internal error: can't make ast::Variable for null type";
return nullptr;
type::Access ParserImpl::VarAccess(const Type* storage_type, type::AddressSpace address_space) {
if (address_space != type::AddressSpace::kStorage) {
return type::Access::kUndefined;
}
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;
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) {
if (storage_type == nullptr) {
Fail() << "internal error: can't make ast::Variable for null type";
return nullptr;
}
// Handle variables (textures and samplers) are always in the handle

24
src/tint/reader/spirv/parser_impl.h
View File

@ -420,10 +420,17 @@ class ParserImpl : Reader {
/// @returns a list of SPIR-V decorations.
DecorationList GetMemberPipelineDecorations(const Struct& struct_type, int member_index);
/// @param storage_type the 'var' storage type
/// @param address_space the 'var' address space
/// @returns the access mode for a 'var' declaration with the given storage type and address
/// space.
type::Access VarAccess(const Type* storage_type, type::AddressSpace address_space);
/// Creates an AST 'var' node for a SPIR-V ID, including any attached decorations, unless it's
/// an ignorable builtin variable.
/// @param id the SPIR-V result ID
/// @param address_space the address space, which cannot be type::AddressSpace::kUndefined
/// @param access the access
/// @param storage_type the storage type of the variable
/// @param initializer the variable initializer
/// @param decorations the variable decorations
@ -431,6 +438,7 @@ class ParserImpl : Reader {
/// in the error case
const ast::Var* MakeVar(uint32_t id,
type::AddressSpace address_space,
type::Access access,
const Type* storage_type,
const ast::Expression* initializer,
AttributeList decorations);
@ -659,13 +667,23 @@ class ParserImpl : Reader {
/// error
const Type* GetHandleTypeForSpirvHandle(const spvtools::opt::Instruction& obj);
/// ModuleVariable describes a module scope variable
struct ModuleVariable {
/// The AST variable node.
const ast::Var* var = nullptr;
/// The address space of the var
type::AddressSpace address_space = type::AddressSpace::kUndefined;
/// The access mode of the var
type::Access access = type::Access::kUndefined;
};
/// Returns the AST variable for the SPIR-V ID of a module-scope variable,
/// or null if there isn't one.
/// @param id a SPIR-V ID
/// @returns the AST variable or null.
const ast::Var* GetModuleVariable(uint32_t id) {
ModuleVariable GetModuleVariable(uint32_t id) {
auto entry = module_variable_.Find(id);
return entry ? *entry : nullptr;
return entry ? *entry : ModuleVariable{};
}
/// Returns the channel component type corresponding to the given image
@ -885,7 +903,7 @@ class ParserImpl : Reader {
std::unordered_map<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.

2
src/tint/reader/wgsl/classify_template_args.cc
View File

@ -70,8 +70,8 @@ void ClassifyTemplateArguments(std::vector<Token>& tokens) {
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) {

75
src/tint/reader/wgsl/parser_impl.cc
View File

@ -204,23 +204,6 @@ ParserImpl::FunctionHeader::~FunctionHeader() = default;
ParserImpl::FunctionHeader& ParserImpl::FunctionHeader::operator=(const FunctionHeader& rhs) =
default;
ParserImpl::VarDeclInfo::VarDeclInfo() = default;
ParserImpl::VarDeclInfo::VarDeclInfo(const VarDeclInfo&) = default;
ParserImpl::VarDeclInfo::VarDeclInfo(Source source_in,
std::string name_in,
type::AddressSpace address_space_in,
type::Access access_in,
ast::Type type_in)
: source(std::move(source_in)),
name(std::move(name_in)),
address_space(address_space_in),
access(access_in),
type(type_in) {}
ParserImpl::VarDeclInfo::~VarDeclInfo() = default;
ParserImpl::ParserImpl(Source::File const* file) : file_(file) {}
ParserImpl::~ParserImpl() = default;
@ -614,8 +597,8 @@ Maybe<const ast::Variable*> ParserImpl::global_variable_decl(AttributeList& attr
TINT_DEFER(attrs.Clear());
return builder_.Var(decl->source, // source
decl->name, // symbol
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
@ -756,37 +739,28 @@ Expect<ParserImpl::TypedIdentifier> ParserImpl::expect_ident_with_type_specifier
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 @@ Expect<ast::Type> ParserImpl::expect_type(std::string_view use) {
return type.value;
}
// address_space
// : 'function'
// | 'private'
// | 'workgroup'
// | 'uniform'
// | 'storage'
//
// Note, we also parse `push_constant` from the experimental extension
Expect<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 @@ Maybe<const ast::VariableDeclStatement*> ParserImpl::variable_statement() {
initializer = initializer_expr.value;
}
auto* var = builder_.Var(decl_source, // source
decl->name, // symbol
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
initializer, // initializer
utils::Empty); // attributes
return create<ast::VariableDeclStatement>(var->source, var);
}
@ -2520,7 +2483,7 @@ Expect<const ast::Expression*> ParserImpl::expect_relational_expression_post_una
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 @@ Expect<const ast::Expression*> ParserImpl::expect_expression() {
if (expr.matched) {
return expr.value;
}
return add_error(t, "expected expression");
return add_error(t, "expected expression for " + std::string(use));
}
Expect<utils::Vector<const ast::Expression*, 3>> ParserImpl::expect_expression_list(
@ -2537,7 +2500,7 @@ Expect<utils::Vector<const ast::Expression*, 3>> ParserImpl::expect_expression_l
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;
}

53
src/tint/reader/wgsl/parser_impl.h
View File

@ -148,10 +148,9 @@ class ParserImpl {
/// Constructor for a successful parse.
/// @param val the result value of the parse
/// @param s the optional source of the value
template <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 @@ class ParserImpl {
/// @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 @@ class ParserImpl {
/// The value of a successful parse.
/// Zero-initialized when there was a parse error.
T value{};
/// Optional source of the value.
Source source;
/// True if there was a error parsing.
bool errored = false;
/// True if there was a error parsing.
@ -268,33 +262,14 @@ class ParserImpl {
/// VarDeclInfo contains the parsed information for variable declaration.
struct VarDeclInfo {
/// Constructor
VarDeclInfo();
/// Copy constructor
/// @param other the VarDeclInfo to copy
VarDeclInfo(const VarDeclInfo& other);
/// Constructor
/// @param source_in variable declaration source
/// @param name_in variable name
/// @param address_space_in variable address space
/// @param access_in variable access control
/// @param type_in variable type
VarDeclInfo(Source source_in,
std::string name_in,
type::AddressSpace address_space_in,
type::Access access_in,
ast::Type type_in);
/// Destructor
~VarDeclInfo();
/// Variable declaration source
Source source;
/// Variable name
std::string name;
/// Variable address space
type::AddressSpace address_space = type::AddressSpace::kUndefined;
const ast::Expression* address_space = nullptr;
/// Variable access control
type::Access access = type::Access::kUndefined;
const ast::Expression* access = nullptr;
/// Variable type
ast::Type type;
};
@ -302,9 +277,9 @@ class ParserImpl {
/// VariableQualifier contains the parsed information for a variable qualifier
struct VariableQualifier {
/// The variable's address space
type::AddressSpace address_space = type::AddressSpace::kUndefined;
const ast::Expression* address_space = nullptr;
/// The variable's access control
type::Access access = type::Access::kUndefined;
const ast::Expression* access = nullptr;
};
/// MatrixDimensions contains the column and row information for a matrix
@ -447,10 +422,6 @@ class ParserImpl {
/// Parses a `type_specifier` grammar element
/// @returns the parsed Type or nullptr if none matched.
Maybe<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 @@ class ParserImpl {
/// 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 @@ class ParserImpl {
/// @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

62
src/tint/reader/wgsl/parser_impl_address_space_test.cc
View File

@ -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

13
src/tint/reader/wgsl/parser_impl_error_msg_test.cc
View File

@ -799,7 +799,7 @@ var x : texture_storage_2d<r32uint, read;
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,18 +1123,9 @@ var i : ptr<private, u32;
)");
}
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
^
)");

30
src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
View File

@ -31,10 +31,8 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithoutInitializer) {
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 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithInitializer) {
ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
EXPECT_EQ(var->declared_address_space, type::AddressSpace::kPrivate);
ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "private");
EXPECT_EQ(var->source.range.begin.line, 1u);
EXPECT_EQ(var->source.range.begin.column, 14u);
@ -83,11 +80,8 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithAttribute) {
ASSERT_NE(var, nullptr);
ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
ASSERT_NE(var->type, nullptr);
ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform);
ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform");
EXPECT_EQ(var->source.range.begin.line, 1u);
EXPECT_EQ(var->source.range.begin.column, 36u);
@ -116,10 +110,8 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithAttribute_MulitpleGroups) {
ASSERT_NE(var, nullptr);
ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
ASSERT_NE(var->type, nullptr);
ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform);
ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform");
EXPECT_EQ(var->source.range.begin.line, 1u);
EXPECT_EQ(var->source.range.begin.column, 36u);
@ -162,19 +154,5 @@ TEST_F(ParserImplTest, GlobalVariableDecl_InvalidConstExpr) {
EXPECT_EQ(p->error(), "1:24: missing initializer for 'var' declaration");
}
TEST_F(ParserImplTest, GlobalVariableDecl_InvalidVariableDecl) {
auto p = parser("var<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

10
src/tint/reader/wgsl/parser_impl_type_decl_test.cc
View File

@ -137,7 +137,7 @@ TEST_P(VecMissingType, Handles_Missing_Type) {
EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error());
ASSERT_EQ(p->error(), "1:6: expected expression");
ASSERT_EQ(p->error(), "1:6: expected expression for type template argument list");
}
INSTANTIATE_TEST_SUITE_P(ParserImplTest,
VecMissingType,
@ -211,7 +211,7 @@ TEST_F(ParserImplTest, TypeDecl_Ptr_MissingAddressSpace) {
EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error());
ASSERT_EQ(p->error(), R"(1:5: expected expression)");
ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)");
}
TEST_F(ParserImplTest, TypeDecl_Ptr_MissingParams) {
@ -221,7 +221,7 @@ TEST_F(ParserImplTest, TypeDecl_Ptr_MissingParams) {
EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error());
ASSERT_EQ(p->error(), R"(1:5: expected expression)");
ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)");
}
TEST_F(ParserImplTest, TypeDecl_Atomic) {
@ -256,7 +256,7 @@ TEST_F(ParserImplTest, TypeDecl_Atomic_MissingType) {
EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error());
ASSERT_EQ(p->error(), "1:8: expected expression");
ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list");
}
TEST_F(ParserImplTest, TypeDecl_Array_AbstractIntLiteralSize) {
@ -431,7 +431,7 @@ TEST_P(MatrixMissingType, Handles_Missing_Type) {
EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error());
ASSERT_EQ(p->error(), "1:8: expected expression");
ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list");
}
INSTANTIATE_TEST_SUITE_P(ParserImplTest,
MatrixMissingType,

18
src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
View File

@ -84,8 +84,7 @@ TEST_F(ParserImplTest, VariableDecl_WithAddressSpace) {
EXPECT_EQ(v->name, "my_var");
ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32");
EXPECT_EQ(v->address_space, type::AddressSpace::kPrivate);
ast::CheckIdentifier(p->builder().Symbols(), v->address_space, "private");
EXPECT_EQ(v->source.range.begin.line, 1u);
EXPECT_EQ(v->source.range.begin.column, 14u);
@ -102,20 +101,7 @@ TEST_F(ParserImplTest, VariableDecl_WithPushConstant) {
EXPECT_EQ(v->name, "my_var");
ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32");
EXPECT_EQ(v->address_space, type::AddressSpace::kPushConstant);
}
TEST_F(ParserImplTest, VariableDecl_InvalidAddressSpace) {
auto p = parser("var<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

37
src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
View File

@ -12,6 +12,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ast/test_helper.h"
#include "src/tint/reader/wgsl/parser_impl_test_helper.h"
namespace tint::reader::wgsl {
@ -31,14 +32,23 @@ class VariableQualifierTest : public ParserImplTestWithParam<VariableStorageData
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 @@ INSTANTIATE_TEST_SUITE_P(
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 @@ TEST_F(ParserImplTest, VariableQualifier_MissingGreaterThan) {
EXPECT_TRUE(p->has_error());
EXPECT_TRUE(sc.errored);
EXPECT_FALSE(sc.matched);
EXPECT_EQ(p->error(), "1:9: expected '>' for variable declaration");
EXPECT_EQ(p->error(), "1:1: missing closing '>' for variable declaration");
}
} // namespace

9
src/tint/resolver/address_space_validation_test.cc
View File

@ -32,8 +32,9 @@ TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_NoAddressSpace_Fail) {
GlobalVar(Source{{12, 34}}, "g", ty.f32());
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"12:34 error: module-scope 'var' declaration must have a address space");
EXPECT_EQ(
r()->error(),
R"(12:34 error: module-scope 'var' declarations that are not of texture or sampler types must provide an address space)");
}
TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NoAddressSpace_Fail) {
@ -469,7 +470,7 @@ TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_NotStorage_AccessMode)
EXPECT_EQ(
r()->error(),
R"(12:34 error: only variables in <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 @@ TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NotStorage_AccessMode) {
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) {

39
src/tint/resolver/dependency_graph.cc
View File

@ -183,13 +183,9 @@ class DependencyScanner {
Declare(func->name->symbol, func);
TraverseFunction(func);
},
[&](const ast::Variable* var) {
Declare(var->name->symbol, var);
TraverseTypeExpression(var->type);
TraverseAttributes(var->attributes);
if (var->initializer) {
TraverseValueExpression(var->initializer);
}
[&](const ast::Variable* v) {
Declare(v->name->symbol, v);
TraverseVariable(v);
},
[&](const ast::DiagnosticDirective*) {
// Diagnostic directives do not affect the dependency graph.
@ -204,8 +200,18 @@ class DependencyScanner {
}
private:
/// Traverses the function, performing symbol resolution and determining
/// global dependencies.
/// Traverses the variable, performing symbol resolution.
void TraverseVariable(const ast::Variable* v) {
if (auto* var = v->As<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 @@ class DependencyScanner {
if (auto* shadows = scope_stack_.Get(v->variable->name->symbol)) {
graph_.shadows.Add(v->variable, shadows);
}
TraverseTypeExpression(v->variable->type);
TraverseValueExpression(v->variable->initializer);
TraverseVariable(v->variable);
Declare(v->variable->name->symbol, v->variable);
},
[&](const ast::WhileStatement* w) {
@ -345,6 +350,18 @@ class DependencyScanner {
TraverseExpression(root, "type", "references");
}
/// Traverses the expression @p root_expr for the intended use as an address space, performing
/// symbol resolution and determining global dependencies.
void TraverseAddressSpaceExpression(const ast::Expression* root) {
TraverseExpression(root, "address space", "references");
}
/// Traverses the expression @p root_expr for the intended use as an access, performing symbol
/// resolution and determining global dependencies.
void TraverseAccessExpression(const ast::Expression* root) {
TraverseExpression(root, "access", "references");
}
/// Traverses the expression @p root_expr for the intended use as a call target, performing
/// symbol resolution and determining global dependencies.
void TraverseCallableExpression(const ast::Expression* root) {

4
src/tint/resolver/dependency_graph_test.cc
View File

@ -1437,13 +1437,13 @@ TEST_P(ResolverDependencyGraphResolveToAddressSpace, Resolve) {
<< resolved->String(Symbols(), Diagnostics());
}
TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalVar) {
TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalConst) {
const auto use = std::get<0>(GetParam());
const auto builtin = std::get<1>(GetParam());
const auto symbol = Symbols().New(utils::ToString(builtin));
SymbolTestHelper helper(this);
auto* decl = helper.Add(SymbolDeclKind::GlobalVar, symbol);
auto* decl = helper.Add(SymbolDeclKind::GlobalConst, symbol);
auto* ident = helper.Add(use, symbol);
helper.Build();

13
src/tint/resolver/function_validation_test.cc
View File

@ -1088,15 +1088,14 @@ TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceWithExtension) {
param.expectation == Expectation::kPassWithFullPtrParameterExtension) {
ASSERT_TRUE(r()->Resolve()) << r()->error();
} else {
std::stringstream ss;
ss << param.address_space;
EXPECT_FALSE(r()->Resolve());
if (param.expectation == Expectation::kInvalid) {
EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" + ss.str() + "'");
EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" +
utils::ToString(param.address_space) + "'");
} else {
EXPECT_EQ(r()->error(),
"12:34 error: function parameter of pointer type cannot be in '" + ss.str() +
"' address space");
"12:34 error: function parameter of pointer type cannot be in '" +
utils::ToString(param.address_space) + "' address space");
}
}
}
@ -1105,8 +1104,8 @@ INSTANTIATE_TEST_SUITE_P(
ResolverFunctionParameterValidationTest,
testing::Values(
TestParams{type::AddressSpace::kUndefined, Expectation::kInvalid},
TestParams{type::AddressSpace::kIn, Expectation::kInvalid},
TestParams{type::AddressSpace::kOut, Expectation::kInvalid},
TestParams{type::AddressSpace::kIn, Expectation::kAlwaysFail},
TestParams{type::AddressSpace::kOut, Expectation::kAlwaysFail},
TestParams{type::AddressSpace::kUniform, Expectation::kPassWithFullPtrParameterExtension},
TestParams{type::AddressSpace::kWorkgroup, Expectation::kPassWithFullPtrParameterExtension},
TestParams{type::AddressSpace::kHandle, Expectation::kInvalid},

20
src/tint/resolver/resolver.cc
View File

@ -473,8 +473,14 @@ sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
return nullptr;
}
auto address_space = var->declared_address_space;
if (address_space == type::AddressSpace::kUndefined) {
auto address_space = type::AddressSpace::kUndefined;
if (var->declared_address_space) {
auto expr = AddressSpaceExpression(var->declared_address_space);
if (!expr) {
return nullptr;
}
address_space = expr->Value();
} else {
// No declared address space. Infer from usage / type.
if (!is_global) {
address_space = type::AddressSpace::kFunction;
@ -494,8 +500,14 @@ sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
return nullptr;
}
auto access = var->declared_access;
if (access == type::Access::kUndefined) {
auto access = type::Access::kUndefined;
if (var->declared_access) {
auto expr = AccessExpression(var->declared_access);
if (!expr) {
return nullptr;
}
access = expr->Value();
} else {
access = DefaultAccessForAddressSpace(address_space);
}

11
src/tint/resolver/validation_test.cc
View File

@ -295,22 +295,23 @@ TEST_F(ResolverValidationTest, AddressSpace_FunctionVariableI32) {
TEST_F(ResolverValidationTest, AddressSpace_SamplerExplicitAddressSpace) {
auto t = ty.sampler(type::SamplerKind::kSampler);
GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a));
GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kPrivate, Binding(0_a), Group(0_a));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
R"(12:34 error: variables of type 'sampler' must not have a address space)");
R"(12:34 error: variables of type 'sampler' must not specifiy an address space)");
}
TEST_F(ResolverValidationTest, AddressSpace_TextureExplicitAddressSpace) {
auto t = ty.sampled_texture(type::TextureDimension::k1d, ty.f32());
GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a));
GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kFunction, Binding(0_a), Group(0_a));
EXPECT_FALSE(r()->Resolve()) << r()->error();
EXPECT_EQ(r()->error(),
R"(12:34 error: variables of type 'texture_1d<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) {

31
src/tint/resolver/validator.cc
View File

@ -299,7 +299,7 @@ bool Validator::Pointer(const ast::TemplatedIdentifier* a, const type::Pointer*
// * For the storage address space, the access mode is optional, and defaults to read.
// * For other address spaces, the access mode must not be written.
if (s->AddressSpace() != type::AddressSpace::kStorage) {
AddError("only pointers in <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 Validator::GlobalVariable(
}
bool ok = Switch(
decl, //
[&](const ast::Var*) {
[&](const ast::Var* var) {
if (auto* init = global->Initializer();
init && init->Stage() > sem::EvaluationStage::kOverride) {
AddError("module-scope 'var' initializer must be a constant or override-expression",
@ -612,8 +612,11 @@ bool Validator::GlobalVariable(
return false;
}
if (global->AddressSpace() == type::AddressSpace::kUndefined) {
AddError("module-scope 'var' declaration must have a address space", decl->source);
if (!var->declared_address_space && !global->Type()->UnwrapRef()->is_handle()) {
AddError(
"module-scope 'var' declarations that are not of texture or sampler types must "
"provide an address space",
decl->source);
return false;
}
@ -696,25 +699,23 @@ bool Validator::Var(const sem::Variable* v) const {
return false;
}
if (store_ty->is_handle()) {
if (var->declared_address_space != type::AddressSpace::kUndefined) {
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.
// 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",
"' 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 @@ bool Validator::Var(const sem::Variable* v) const {
}
if (IsValidationEnabled(var->attributes, ast::DisabledValidation::kIgnoreAddressSpace) &&
(var->declared_address_space == type::AddressSpace::kIn ||
var->declared_address_space == type::AddressSpace::kOut)) {
(v->AddressSpace() == type::AddressSpace::kIn ||
v->AddressSpace() == type::AddressSpace::kOut)) {
AddError("invalid use of input/output address space", var->source);
return false;
}

22
src/tint/transform/binding_remapper.cc
View File

@ -122,11 +122,12 @@ Transform::ApplyResult BindingRemapper::Apply(const Program* src,
// Replace any access controls.
auto ac_it = remappings->access_controls.find(from);
if (ac_it != remappings->access_controls.end()) {
type::Access ac = ac_it->second;
if (ac == type::Access::kUndefined) {
b.Diagnostics().add_error(
diag::System::Transform,
"invalid access mode (" + std::to_string(static_cast<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 @@ Transform::ApplyResult BindingRemapper::Apply(const Program* src,
}
auto* ty = sem->Type()->UnwrapRef();
auto inner_ty = CreateASTTypeFor(ctx, ty);
auto* new_var = b.Var(ctx.Clone(var->source), ctx.Clone(var->name->symbol),
inner_ty, var->declared_address_space, ac,
ctx.Clone(var->initializer), ctx.Clone(var->attributes));
auto* new_var =
b.create<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);
}

170
src/tint/transform/canonicalize_entry_point_io_test.cc
View File

@ -67,11 +67,11 @@ fn frag_main(@location(1) loc1 : f32,
)";
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 @@ fn frag_main(@location(0) loc0 : f32,
)";
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 @@ struct FragLocations {
)";
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 @@ fn frag_main() -> @builtin(frag_depth) f32 {
)";
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 @@ fn frag_main() -> FragOutput {
)";
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 @@ struct FragOutput {
)";
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 @@ fn frag_main2(inputs : FragmentInput) {
)";
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 @@ struct FragmentInput {
)";
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 @@ fn frag_main(inputs : FragmentInterface) -> FragmentInterface {
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 @@ struct FragmentInterface {
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 @@ fn vert_main() -> @builtin(position) vec4<f32> {
)";
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 @@ fn vert_main() -> VertOut {
)";
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 @@ struct VertOut {
)";
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 @@ fn vert_main(collide : VertIn1, collide_1 : VertIn2) -> VertOut {
)";
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 @@ struct VertOut {
)";
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 @@ fn main(@builtin(sample_index) sample_index : u32,
)";
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 @@ fn fragment_main(@builtin(sample_index) sample_index : u32,
)";
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 @@ fn vertex_main(@builtin(vertex_index) vertexID : u32,
)";
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)));

5
src/tint/transform/clamp_frag_depth.cc
View File

@ -88,10 +88,11 @@ Transform::ApplyResult ClampFragDepth::Apply(const Program* src, const DataMap&,
// Abort on any use of push constants in the module.
for (auto* global : src->AST().GlobalVariables()) {
if (auto* var = global->As<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));
}
}

10
src/tint/transform/module_scope_var_to_entry_point_param.cc
View File

@ -117,7 +117,6 @@ struct ModuleScopeVarToEntryPointParam::State {
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 @@ struct ModuleScopeVarToEntryPointParam::State {
is_wrapped = true;
}
param_type = ctx.dst->ty.pointer(param_type, sc, var_ast->declared_access);
param_type = sc == type::AddressSpace::kStorage
? ctx.dst->ty.pointer(param_type, sc, var->Access())
: ctx.dst->ty.pointer(param_type, sc);
auto* param = ctx.dst->Param(new_var_symbol, param_type, attributes);
ctx.InsertFront(func->params, param);
is_pointer = true;
@ -228,7 +229,6 @@ struct ModuleScopeVarToEntryPointParam::State {
const sem::Variable* var,
Symbol new_var_symbol,
bool& is_pointer) {
auto* var_ast = var->Declaration()->As<ast::Var>();
auto* ty = var->Type()->UnwrapRef();
auto param_type = CreateASTTypeFor(ctx, ty);
auto sc = var->AddressSpace();
@ -254,7 +254,9 @@ struct ModuleScopeVarToEntryPointParam::State {
// Use a pointer for non-handle types.
utils::Vector<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.

36
src/tint/transform/module_scope_var_to_entry_point_param_test.cc
View File

@ -435,7 +435,7 @@ struct S {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct S {
auto* expect = R"(
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct tint_symbol_1 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct tint_symbol_1 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct tint_symbol_2 {
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 @@ struct tint_symbol_2 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<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 @@ struct tint_symbol_1 {
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 @@ struct tint_symbol_1 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct tint_symbol_1 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct tint_symbol_1 {
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<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 @@ struct S {
fn no_uses() {
}
fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<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 @@ fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(di
}
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<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 @@ var<storage> s : S;
auto* expect = R"(
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<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 foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(di
fn no_uses() {
}
fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<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);
}

12
src/tint/transform/renamer_test.cc
View File

@ -1927,8 +1927,10 @@ std::vector<const char*> Identifiers() {
out.push_back(ident);
}
for (auto* ident : type::kAddressSpaceStrings) {
if (!utils::HasPrefix(ident, "_")) {
out.push_back(ident);
}
}
for (auto* ident : type::kTexelFormatStrings) {
out.push_back(ident);
}
@ -1940,24 +1942,24 @@ std::vector<const char*> Identifiers() {
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;
}
)");

2
src/tint/transform/std140.cc
View File

@ -350,8 +350,8 @@ struct Std140::State {
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);
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);

3
src/tint/transform/zero_init_workgroup_memory.cc
View File

@ -36,7 +36,8 @@ namespace {
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;
}
}

14
src/tint/type/address_space.cc
View File

@ -28,6 +28,12 @@ namespace tint::type {
/// @param str the string to parse
/// @returns the parsed enum, or AddressSpace::kUndefined if the string could not be parsed.
AddressSpace ParseAddressSpace(std::string_view str) {
if (str == "__in") {
return AddressSpace::kIn;
}
if (str == "__out") {
return AddressSpace::kOut;
}
if (str == "function") {
return AddressSpace::kFunction;
}
@ -53,14 +59,14 @@ std::ostream& operator<<(std::ostream& out, AddressSpace value) {
switch (value) {
case AddressSpace::kUndefined:
return out << "undefined";
case AddressSpace::kIn:
return out << "__in";
case AddressSpace::kOut:
return out << "__out";
case AddressSpace::kFunction:
return out << "function";
case AddressSpace::kHandle:
return out << "handle";
case AddressSpace::kIn:
return out << "in";
case AddressSpace::kOut:
return out << "out";
case AddressSpace::kPrivate:
return out << "private";
case AddressSpace::kPushConstant:

6
src/tint/type/address_space.h
View File

@ -30,10 +30,10 @@ namespace tint::type {
/// Address space of a given pointer.
enum class AddressSpace {
kUndefined,
kIn,
kOut,
kFunction,
kHandle, // Tint-internal enum entry - not parsed
kIn, // Tint-internal enum entry - not parsed
kOut, // Tint-internal enum entry - not parsed
kPrivate,
kPushConstant,
kStorage,
@ -52,7 +52,7 @@ std::ostream& operator<<(std::ostream& out, AddressSpace value);
AddressSpace ParseAddressSpace(std::string_view str);
constexpr const char* kAddressSpaceStrings[] = {
"function", "private", "push_constant", "storage", "uniform", "workgroup",
"__in", "__out", "function", "private", "push_constant", "storage", "uniform", "workgroup",
};
/// @returns true if the AddressSpace is host-shareable

65
src/tint/type/address_space_bench.cc
View File

@ -31,15 +31,62 @@ namespace {
void AddressSpaceParser(::benchmark::State& state) {
const char* kStrings[] = {
"fccnctin", "ucti3", "functVon", "function", "1unction",
"unJtqqon", "llun77tion", "ppqqivtHH", "prcv", "bivaGe",
"private", "priviive", "8WWivate", "pxxvate", "pXh_cggnstant",
"pX_Vonstanu", "push_consta3t", "push_constant", "push_constanE", "push_TTPnstant",
"puxxdh_constan", "s44orage", "stSSraVVe", "RtoR22e", "storage",
"sFra9e", "stoage", "VOORRHge", "unfoym", "llnnrrf77rm",
"unif4r00", "uniform", "nfoom", "zzform", "uiiippo1",
"workgrouXX", "wor55gro99nII", "wrrrkgroSSaHH", "workgroup", "kkrHoup",
"jgkrouRR", "wokroub",
"ccin",
"3",
"_Vin",
"__in",
"1_in",
"_qiJ",
"_lli77",
"__qHupp",
"vt",
"G_bt",
"__out",
"__viut",
"__8WWt",
"Mxxou",
"fuXggton",
"fuXtou",
"funct3on",
"function",
"funEtion",
"PPncTTion",
"xxuncddon",
"p44ivate",
"prSSvaVVe",
"RriR22e",
"private",
"pFva9e",
"priate",
"VOORRHte",
"push_constyn",
"punnh_crr77stallt",
"pu4h_cons00ant",
"push_constant",
"puoo_costan",
"ushzzcnstant",
"push_coii11apt",
"storaXXe",
"9II5tnnrage",
"stoaSSrHHYe",
"storage",
"stkke",
"jtogRa",
"sbrag",
"unifojm",
"niform",
"qform",
"uniform",
"uniNNrm",
"nifrvv",
"QQiform",
"workrorf",
"workjroup",
"wNNorkrou2",
"workgroup",
"workgrop",
"rrorkgroup",
"workgroGp",
};
for (auto _ : state) {
for (auto* str : kStrings) {

23
src/tint/type/address_space_test.cc
View File

@ -42,6 +42,8 @@ inline std::ostream& operator<<(std::ostream& out, Case c) {
}
static constexpr Case kValidCases[] = {
{"__in", AddressSpace::kIn},
{"__out", AddressSpace::kOut},
{"function", AddressSpace::kFunction},
{"private", AddressSpace::kPrivate},
{"push_constant", AddressSpace::kPushConstant},
@ -51,15 +53,18 @@ static constexpr Case kValidCases[] = {
};
static constexpr Case kInvalidCases[] = {
{"fccnctin", AddressSpace::kUndefined}, {"ucti3", AddressSpace::kUndefined},
{"functVon", AddressSpace::kUndefined}, {"priv1te", AddressSpace::kUndefined},
{"pqiJate", AddressSpace::kUndefined}, {"privat7ll", AddressSpace::kUndefined},
{"pqqsh_pponstHnt", AddressSpace::kUndefined}, {"pus_cnstat", AddressSpace::kUndefined},
{"bus_Gonstant", AddressSpace::kUndefined}, {"storiive", AddressSpace::kUndefined},
{"8WWorage", AddressSpace::kUndefined}, {"sxxrage", AddressSpace::kUndefined},
{"uXforgg", AddressSpace::kUndefined}, {"nfoXm", AddressSpace::kUndefined},
{"unif3rm", AddressSpace::kUndefined}, {"workgroEp", AddressSpace::kUndefined},
{"woTTPkroup", AddressSpace::kUndefined}, {"ddorkroxxp", AddressSpace::kUndefined},
{"ccin", AddressSpace::kUndefined}, {"3", AddressSpace::kUndefined},
{"_Vin", AddressSpace::kUndefined}, {"__ou1", AddressSpace::kUndefined},
{"qq_Jt", AddressSpace::kUndefined}, {"__oll7t", AddressSpace::kUndefined},
{"qquntppHon", AddressSpace::kUndefined}, {"cnciv", AddressSpace::kUndefined},
{"funGion", AddressSpace::kUndefined}, {"priviive", AddressSpace::kUndefined},
{"8WWivate", AddressSpace::kUndefined}, {"pxxvate", AddressSpace::kUndefined},
{"pXh_cggnstant", AddressSpace::kUndefined}, {"pX_Vonstanu", AddressSpace::kUndefined},
{"push_consta3t", AddressSpace::kUndefined}, {"Etorage", AddressSpace::kUndefined},
{"sPTTrage", AddressSpace::kUndefined}, {"storadxx", AddressSpace::kUndefined},
{"u44iform", AddressSpace::kUndefined}, {"unSSfoVVm", AddressSpace::kUndefined},
{"RniR22m", AddressSpace::kUndefined}, {"w9rFroup", AddressSpace::kUndefined},
{"workgoup", AddressSpace::kUndefined}, {"woVROOrHup", AddressSpace::kUndefined},
};
using AddressSpaceParseTest = testing::TestWithParam<Case>;

35
src/tint/writer/wgsl/generator_impl.cc
View File

@ -34,8 +34,6 @@
#include "src/tint/ast/workgroup_attribute.h"
#include "src/tint/sem/struct.h"
#include "src/tint/sem/switch_statement.h"
#include "src/tint/type/access.h"
#include "src/tint/type/texture_dimension.h"
#include "src/tint/utils/math.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/writer/float_to_string.h"
@ -376,24 +374,6 @@ bool GeneratorImpl::EmitImageFormat(std::ostream& out, const type::TexelFormat f
return true;
}
bool GeneratorImpl::EmitAccess(std::ostream& out, const type::Access access) {
switch (access) {
case type::Access::kRead:
out << "read";
return true;
case type::Access::kWrite:
out << "write";
return true;
case type::Access::kReadWrite:
out << "read_write";
return true;
default:
break;
}
diagnostics_.add_error(diag::System::Writer, "unknown access");
return false;
}
bool GeneratorImpl::EmitStructType(const ast::Struct* str) {
if (str->attributes.Length()) {
if (!EmitAttributes(line(), str->attributes)) {
@ -473,17 +453,18 @@ bool GeneratorImpl::EmitVariable(std::ostream& out, const ast::Variable* v) {
v, //
[&](const ast::Var* var) {
out << "var";
auto address_space = var->declared_address_space;
auto ac = var->declared_access;
if (address_space != type::AddressSpace::kUndefined || ac != type::Access::kUndefined) {
out << "<" << address_space;
if (ac != type::Access::kUndefined) {
if (var->declared_address_space || var->declared_access) {
out << "<";
TINT_DEFER(out << ">");
if (!EmitExpression(out, var->declared_address_space)) {
return false;
}
if (var->declared_access) {
out << ", ";
if (!EmitAccess(out, ac)) {
if (!EmitExpression(out, var->declared_access)) {
return false;
}
}
out << ">";
}
return true;
},

6
src/tint/writer/wgsl/generator_impl.h
View File

@ -35,7 +35,6 @@
#include "src/tint/ast/unary_op_expression.h"
#include "src/tint/program.h"
#include "src/tint/sem/struct.h"
#include "src/tint/type/storage_texture.h"
#include "src/tint/writer/text_generator.h"
namespace tint::writer::wgsl {
@ -209,11 +208,6 @@ class GeneratorImpl : public TextGenerator {
/// @param fmt the format to generate
/// @returns true if the format is emitted
bool EmitImageFormat(std::ostream& out, const type::TexelFormat fmt);
/// Handles emitting an access control
/// @param out the output stream
/// @param access the access to generate
/// @returns true if the access is emitted
bool EmitAccess(std::ostream& out, const type::Access access);
/// Handles a unary op expression
/// @param out the output stream
/// @param expr the expression to emit