tint: Fix HLSL emission for out-of-order storage / uniform buffers
Recent changes to DecomposeMemoryAccess meant we lost the dependency information between the user of a module-scope variable of the storage / uniform address space and the variable.
Add dependency information to ast::InternalAttribute so this can be tracked.
This change also means that symbol renaming after the DecomposeMemoryAccess should work.
Fixed: tint:1860
Change-Id: Icfa2925f95c2ac50702522df514cd11bde727546
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/122660
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/ast/disable_validation_attribute.cc b/src/tint/ast/disable_validation_attribute.cc
index eff1c1f..27eaef0 100644
--- a/src/tint/ast/disable_validation_attribute.cc
+++ b/src/tint/ast/disable_validation_attribute.cc
@@ -23,7 +23,7 @@
DisableValidationAttribute::DisableValidationAttribute(ProgramID pid,
NodeID nid,
DisabledValidation val)
- : Base(pid, nid), validation(val) {}
+ : Base(pid, nid, utils::Empty), validation(val) {}
DisableValidationAttribute::~DisableValidationAttribute() = default;
diff --git a/src/tint/ast/internal_attribute.cc b/src/tint/ast/internal_attribute.cc
index 1b4ca9e..c5c5f2e 100644
--- a/src/tint/ast/internal_attribute.cc
+++ b/src/tint/ast/internal_attribute.cc
@@ -14,11 +14,16 @@
#include "src/tint/ast/internal_attribute.h"
+#include <utility>
+
TINT_INSTANTIATE_TYPEINFO(tint::ast::InternalAttribute);
namespace tint::ast {
-InternalAttribute::InternalAttribute(ProgramID pid, NodeID nid) : Base(pid, nid, Source{}) {}
+InternalAttribute::InternalAttribute(ProgramID pid,
+ NodeID nid,
+ utils::VectorRef<const IdentifierExpression*> deps)
+ : Base(pid, nid, Source{}), dependencies(std::move(deps)) {}
InternalAttribute::~InternalAttribute() = default;
diff --git a/src/tint/ast/internal_attribute.h b/src/tint/ast/internal_attribute.h
index 9904af8..36f2a98 100644
--- a/src/tint/ast/internal_attribute.h
+++ b/src/tint/ast/internal_attribute.h
@@ -18,6 +18,12 @@
#include <string>
#include "src/tint/ast/attribute.h"
+#include "src/tint/utils/vector.h"
+
+// Forward declarations
+namespace tint::ast {
+class IdentifierExpression;
+} // namespace tint::ast
namespace tint::ast {
@@ -29,7 +35,10 @@
/// Constructor
/// @param program_id the identifier of the program that owns this node
/// @param nid the unique node identifier
- explicit InternalAttribute(ProgramID program_id, NodeID nid);
+ /// @param deps a list of identifiers that this attribute is dependent on
+ InternalAttribute(ProgramID program_id,
+ NodeID nid,
+ utils::VectorRef<const IdentifierExpression*> deps);
/// Destructor
~InternalAttribute() override;
@@ -40,6 +49,9 @@
/// @returns the WGSL name for the attribute
std::string Name() const override;
+
+ /// A list of identifiers that this attribute is dependent on
+ const utils::Vector<const IdentifierExpression*, 1> dependencies;
};
} // namespace tint::ast
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index b0a863d..5c09c2a 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -1977,4 +1977,40 @@
} // namespace
} // namespace MustUseTests
+namespace InternalAttributeDeps {
+namespace {
+
+class TestAttribute : public Castable<TestAttribute, ast::InternalAttribute> {
+ public:
+ TestAttribute(ProgramID pid, ast::NodeID nid, const ast::IdentifierExpression* dep)
+ : Base(pid, nid, utils::Vector{dep}) {}
+ std::string InternalName() const override { return "test_attribute"; }
+ const Cloneable* Clone(CloneContext*) const override { return nullptr; }
+};
+
+using InternalAttributeDepsTest = ResolverTest;
+TEST_F(InternalAttributeDepsTest, Dependency) {
+ auto* ident = Expr("v");
+ auto* attr = ASTNodes().Create<TestAttribute>(ID(), AllocateNodeID(), ident);
+ auto* f = Func("f", utils::Empty, ty.void_(), utils::Empty, utils::Vector{attr});
+ auto* v = GlobalVar("v", ty.i32(), builtin::AddressSpace::kPrivate);
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* user = As<sem::VariableUser>(Sem().Get(ident));
+ ASSERT_NE(user, nullptr);
+
+ auto* var = Sem().Get(v);
+ EXPECT_EQ(user->Variable(), var);
+
+ auto* fn = Sem().Get(f);
+ EXPECT_THAT(fn->DirectlyReferencedGlobals(), testing::ElementsAre(var));
+ EXPECT_THAT(fn->TransitivelyReferencedGlobals(), testing::ElementsAre(var));
+}
+
+} // namespace
+} // namespace InternalAttributeDeps
+
} // namespace tint::resolver
+
+TINT_INSTANTIATE_TYPEINFO(tint::resolver::InternalAttributeDeps::TestAttribute);
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 8292adb..f718a46 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -414,12 +414,18 @@
TraverseExpression(wg->y);
TraverseExpression(wg->z);
return true;
+ },
+ [&](const ast::InternalAttribute* i) {
+ for (auto* dep : i->dependencies) {
+ TraverseExpression(dep);
+ }
+ return true;
});
if (handled) {
return;
}
- if (attr->IsAnyOf<ast::BuiltinAttribute, ast::DiagnosticAttribute, ast::InternalAttribute,
+ if (attr->IsAnyOf<ast::BuiltinAttribute, ast::DiagnosticAttribute,
ast::InterpolateAttribute, ast::InvariantAttribute, ast::MustUseAttribute,
ast::StageAttribute, ast::StrideAttribute,
ast::StructMemberOffsetAttribute>()) {
diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc
index 0d2cc10..e8fe271 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/src/tint/resolver/intrinsic_table.cc
@@ -1495,7 +1495,7 @@
// Was this overload a constructor or conversion?
if (match.overload->flags.Contains(OverloadFlag::kIsConstructor)) {
- utils::Vector<const sem::Parameter*, 8> params;
+ utils::Vector<sem::Parameter*, 8> params;
params.Reserve(match.parameters.Length());
for (auto& p : match.parameters) {
params.Push(builder.create<sem::Parameter>(
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index cd426f7..f51e5cc 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -856,9 +856,9 @@
sem::Function* Resolver::Function(const ast::Function* decl) {
Mark(decl->name);
- uint32_t parameter_index = 0;
- utils::Hashmap<Symbol, Source, 8> parameter_names;
- utils::Vector<sem::Parameter*, 8> parameters;
+ auto* func = builder_->create<sem::Function>(decl);
+ builder_->Sem().Add(decl, func);
+ TINT_SCOPED_ASSIGNMENT(current_function_, func);
validator_.DiagnosticFilters().Push();
TINT_DEFER(validator_.DiagnosticFilters().Pop());
@@ -872,6 +872,8 @@
}
// Resolve all the parameters
+ uint32_t parameter_index = 0;
+ utils::Hashmap<Symbol, Source, 8> parameter_names;
for (auto* param : decl->params) {
Mark(param);
@@ -893,7 +895,7 @@
return nullptr;
}
- parameters.Push(p);
+ func->AddParameter(p);
auto* p_ty = const_cast<type::Type*>(p->Type());
if (auto* str = p_ty->As<sem::Struct>()) {
@@ -923,9 +925,9 @@
} else {
return_type = builder_->create<type::Void>();
}
+ func->SetReturnType(return_type);
// Determine if the return type has a location
- std::optional<uint32_t> return_location;
for (auto* attr : decl->return_type_attributes) {
if (!Attribute(attr)) {
return nullptr;
@@ -936,7 +938,7 @@
if (!value) {
return nullptr;
}
- return_location = value.Get();
+ func->SetReturnLocation(value.Get());
}
}
@@ -963,12 +965,7 @@
}
}
- auto* func =
- builder_->create<sem::Function>(decl, return_type, return_location, std::move(parameters));
ApplyDiagnosticSeverities(func);
- builder_->Sem().Add(decl, func);
-
- TINT_SCOPED_ASSIGNMENT(current_function_, func);
if (!WorkgroupSize(decl)) {
return nullptr;
@@ -2089,7 +2086,7 @@
auto* call_target = struct_ctors_.GetOrCreate(
StructConstructorSig{{str, args.Length(), args_stage}},
[&]() -> sem::ValueConstructor* {
- utils::Vector<const sem::Parameter*, 8> params;
+ utils::Vector<sem::Parameter*, 8> params;
params.Resize(std::min(args.Length(), str->Members().Length()));
for (size_t i = 0, n = params.Length(); i < n; i++) {
params[i] = builder_->create<sem::Parameter>(
@@ -3436,6 +3433,7 @@
[&](const ast::BuiltinAttribute* b) { return BuiltinAttribute(b); },
[&](const ast::DiagnosticAttribute* d) { return DiagnosticControl(d->control); },
[&](const ast::InterpolateAttribute* i) { return InterpolateAttribute(i); },
+ [&](const ast::InternalAttribute* i) { return InternalAttribute(i); },
[&](Default) { return true; });
}
@@ -3460,6 +3458,15 @@
return true;
}
+bool Resolver::InternalAttribute(const ast::InternalAttribute* attr) {
+ for (auto* dep : attr->dependencies) {
+ if (!Expression(dep)) {
+ return false;
+ }
+ }
+ return true;
+}
+
bool Resolver::DiagnosticControl(const ast::DiagnosticControl& control) {
Mark(control.rule_name);
diff --git a/src/tint/resolver/resolver.h b/src/tint/resolver/resolver.h
index fff7506..8acf602 100644
--- a/src/tint/resolver/resolver.h
+++ b/src/tint/resolver/resolver.h
@@ -321,6 +321,10 @@
/// @returns true on success, false on failure
bool InterpolateAttribute(const ast::InterpolateAttribute* attr);
+ /// Resolves the internal attribute @p attr
+ /// @returns true on success, false on failure
+ bool InternalAttribute(const ast::InternalAttribute* attr);
+
/// @param control the diagnostic control
/// @returns true on success, false on failure
bool DiagnosticControl(const ast::DiagnosticControl& control);
diff --git a/src/tint/sem/builtin.cc b/src/tint/sem/builtin.cc
index 8bcd8af..4fd2c34 100644
--- a/src/tint/sem/builtin.cc
+++ b/src/tint/sem/builtin.cc
@@ -25,17 +25,6 @@
TINT_INSTANTIATE_TYPEINFO(tint::sem::Builtin);
namespace tint::sem {
-namespace {
-
-utils::VectorRef<const Parameter*> SetOwner(utils::VectorRef<Parameter*> parameters,
- const tint::sem::CallTarget* owner) {
- for (auto* parameter : parameters) {
- parameter->SetOwner(owner);
- }
- return parameters;
-}
-
-} // namespace
const char* Builtin::str() const {
return sem::str(type_);
@@ -112,7 +101,7 @@
PipelineStageSet supported_stages,
bool is_deprecated,
bool must_use)
- : Base(return_type, SetOwner(std::move(parameters), this), eval_stage, must_use),
+ : Base(return_type, std::move(parameters), eval_stage, must_use),
type_(type),
supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {}
diff --git a/src/tint/sem/call_target.cc b/src/tint/sem/call_target.cc
index 76c1ab1..9ef0059 100644
--- a/src/tint/sem/call_target.cc
+++ b/src/tint/sem/call_target.cc
@@ -23,17 +23,25 @@
namespace tint::sem {
+CallTarget::CallTarget(EvaluationStage stage, bool must_use) : stage_(stage), must_use_(must_use) {}
+
CallTarget::CallTarget(const type::Type* return_type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage,
bool must_use)
- : signature_{return_type, std::move(parameters)}, stage_(stage), must_use_(must_use) {
+ : stage_(stage), must_use_(must_use) {
+ SetReturnType(return_type);
+ for (auto* param : parameters) {
+ AddParameter(param);
+ }
TINT_ASSERT(Semantic, return_type);
}
CallTarget::CallTarget(const CallTarget&) = default;
CallTarget::~CallTarget() = default;
+CallTargetSignature::CallTargetSignature() = default;
+
CallTargetSignature::CallTargetSignature(const type::Type* ret_ty,
utils::VectorRef<const sem::Parameter*> params)
: return_type(ret_ty), parameters(std::move(params)) {}
diff --git a/src/tint/sem/call_target.h b/src/tint/sem/call_target.h
index 4f0ab33..096aa8f 100644
--- a/src/tint/sem/call_target.h
+++ b/src/tint/sem/call_target.h
@@ -28,6 +28,9 @@
/// CallTargetSignature holds the return type and parameters for a call target
struct CallTargetSignature {
/// Constructor
+ CallTargetSignature();
+
+ /// Constructor
/// @param ret_ty the call target return type
/// @param params the call target parameters
CallTargetSignature(const type::Type* ret_ty, utils::VectorRef<const Parameter*> params);
@@ -39,9 +42,9 @@
~CallTargetSignature();
/// The type of the call target return value
- const type::Type* const return_type = nullptr;
+ const type::Type* return_type = nullptr;
/// The parameters of the call target
- const utils::Vector<const sem::Parameter*, 8> parameters;
+ utils::Vector<const sem::Parameter*, 8> parameters;
/// Equality operator
/// @param other the signature to compare this to
@@ -67,13 +70,19 @@
class CallTarget : public Castable<CallTarget, Node> {
public:
/// Constructor
+ /// @param stage the earliest evaluation stage for a call to this target
+ /// @param must_use the result of the call target must be used, i.e. it cannot be used as a call
+ /// statement.
+ CallTarget(EvaluationStage stage, bool must_use);
+
+ /// Constructor
/// @param return_type the return type of the call target
/// @param parameters the parameters for the call target
/// @param stage the earliest evaluation stage for a call to this target
/// @param must_use the result of the call target must be used, i.e. it cannot be used as a call
/// statement.
CallTarget(const type::Type* return_type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage,
bool must_use);
@@ -83,9 +92,20 @@
/// Destructor
~CallTarget() override;
+ /// Sets the call target's return type
+ /// @param ty the parameter
+ void SetReturnType(const type::Type* ty) { signature_.return_type = ty; }
+
/// @return the return type of the call target
const type::Type* ReturnType() const { return signature_.return_type; }
+ /// Adds a parameter to the call target
+ /// @param parameter the parameter
+ void AddParameter(Parameter* parameter) {
+ parameter->SetOwner(this);
+ signature_.parameters.Push(parameter);
+ }
+
/// @return the parameters of the call target
auto& Parameters() const { return signature_.parameters; }
diff --git a/src/tint/sem/function.cc b/src/tint/sem/function.cc
index 0859dd8..ec51b7b 100644
--- a/src/tint/sem/function.cc
+++ b/src/tint/sem/function.cc
@@ -28,29 +28,12 @@
TINT_INSTANTIATE_TYPEINFO(tint::sem::Function);
namespace tint::sem {
-namespace {
-utils::VectorRef<const Parameter*> SetOwner(utils::VectorRef<Parameter*> parameters,
- const tint::sem::CallTarget* owner) {
- for (auto* parameter : parameters) {
- parameter->SetOwner(owner);
- }
- return parameters;
-}
-
-} // namespace
-
-Function::Function(const ast::Function* declaration,
- type::Type* return_type,
- std::optional<uint32_t> return_location,
- utils::VectorRef<Parameter*> parameters)
- : Base(return_type,
- SetOwner(std::move(parameters), this),
- EvaluationStage::kRuntime,
+Function::Function(const ast::Function* declaration)
+ : Base(EvaluationStage::kRuntime,
ast::HasAttribute<ast::MustUseAttribute>(declaration->attributes)),
declaration_(declaration),
- workgroup_size_{1, 1, 1},
- return_location_(return_location) {}
+ workgroup_size_{1, 1, 1} {}
Function::~Function() = default;
diff --git a/src/tint/sem/function.h b/src/tint/sem/function.h
index 528ebef..b157e7d 100644
--- a/src/tint/sem/function.h
+++ b/src/tint/sem/function.h
@@ -54,17 +54,15 @@
/// Constructor
/// @param declaration the ast::Function
- /// @param return_type the return type of the function
- /// @param return_location the location value for the return, if provided
- /// @param parameters the parameters to the function
- Function(const ast::Function* declaration,
- type::Type* return_type,
- std::optional<uint32_t> return_location,
- utils::VectorRef<Parameter*> parameters);
+ explicit Function(const ast::Function* declaration);
/// Destructor
~Function() override;
+ /// Sets the function's return location
+ /// @param return_location the location value
+ void SetReturnLocation(uint32_t return_location) { return_location_ = return_location; }
+
/// @returns the ast::Function declaration
const ast::Function* Declaration() const { return declaration_; }
diff --git a/src/tint/sem/value_constructor.cc b/src/tint/sem/value_constructor.cc
index 62be478..b270414 100644
--- a/src/tint/sem/value_constructor.cc
+++ b/src/tint/sem/value_constructor.cc
@@ -21,7 +21,7 @@
namespace tint::sem {
ValueConstructor::ValueConstructor(const type::Type* type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage)
: Base(type, std::move(parameters), stage, /* must_use */ true) {}
diff --git a/src/tint/sem/value_constructor.h b/src/tint/sem/value_constructor.h
index aef827b..34c3b43 100644
--- a/src/tint/sem/value_constructor.h
+++ b/src/tint/sem/value_constructor.h
@@ -28,7 +28,7 @@
/// @param parameters the constructor parameters
/// @param stage the earliest evaluation stage for the expression
ValueConstructor(const type::Type* type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage);
/// Destructor
diff --git a/src/tint/sem/value_conversion.cc b/src/tint/sem/value_conversion.cc
index 37331af..95587fe 100644
--- a/src/tint/sem/value_conversion.cc
+++ b/src/tint/sem/value_conversion.cc
@@ -19,9 +19,9 @@
namespace tint::sem {
ValueConversion::ValueConversion(const type::Type* type,
- const sem::Parameter* parameter,
+ sem::Parameter* parameter,
EvaluationStage stage)
- : Base(type, utils::Vector<const sem::Parameter*, 1>{parameter}, stage, /* must_use */ true) {}
+ : Base(type, utils::Vector<sem::Parameter*, 1>{parameter}, stage, /* must_use */ true) {}
ValueConversion::~ValueConversion() = default;
diff --git a/src/tint/sem/value_conversion.h b/src/tint/sem/value_conversion.h
index b79caa3..2d2ab38 100644
--- a/src/tint/sem/value_conversion.h
+++ b/src/tint/sem/value_conversion.h
@@ -26,7 +26,7 @@
/// @param type the target type of the cast
/// @param parameter the type cast parameter
/// @param stage the earliest evaluation stage for the expression
- ValueConversion(const type::Type* type, const sem::Parameter* parameter, EvaluationStage stage);
+ ValueConversion(const type::Type* type, sem::Parameter* parameter, EvaluationStage stage);
/// Destructor
~ValueConversion() override;
diff --git a/src/tint/transform/add_block_attribute.cc b/src/tint/transform/add_block_attribute.cc
index c63bd04..c486936 100644
--- a/src/tint/transform/add_block_attribute.cc
+++ b/src/tint/transform/add_block_attribute.cc
@@ -102,7 +102,7 @@
}
AddBlockAttribute::BlockAttribute::BlockAttribute(ProgramID pid, ast::NodeID nid)
- : Base(pid, nid) {}
+ : Base(pid, nid, utils::Empty) {}
AddBlockAttribute::BlockAttribute::~BlockAttribute() = default;
std::string AddBlockAttribute::BlockAttribute::InternalName() const {
return "block";
diff --git a/src/tint/transform/calculate_array_length.cc b/src/tint/transform/calculate_array_length.cc
index a25c89e..9f5b659 100644
--- a/src/tint/transform/calculate_array_length.cc
+++ b/src/tint/transform/calculate_array_length.cc
@@ -71,7 +71,7 @@
} // namespace
CalculateArrayLength::BufferSizeIntrinsic::BufferSizeIntrinsic(ProgramID pid, ast::NodeID nid)
- : Base(pid, nid) {}
+ : Base(pid, nid, utils::Empty) {}
CalculateArrayLength::BufferSizeIntrinsic::~BufferSizeIntrinsic() = default;
std::string CalculateArrayLength::BufferSizeIntrinsic::InternalName() const {
return "intrinsic_buffer_size";
diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc
index 0030ca1..f3d92c5 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -228,7 +228,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), DecomposeMemoryAccess::Intrinsic::Op::kLoad, type,
- address_space, buffer);
+ address_space, builder->Expr(buffer));
}
/// @returns a DecomposeMemoryAccess::Intrinsic attribute that can be applied to a stub function to
@@ -242,7 +242,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), DecomposeMemoryAccess::Intrinsic::Op::kStore,
- type, builtin::AddressSpace::kStorage, buffer);
+ type, builtin::AddressSpace::kStorage, builder->Expr(buffer));
}
/// @returns a DecomposeMemoryAccess::Intrinsic attribute that can be applied to a stub function for
@@ -299,7 +299,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), op, type, builtin::AddressSpace::kStorage,
- buffer);
+ builder->Expr(buffer));
}
/// BufferAccess describes a single storage or uniform buffer access
@@ -692,8 +692,8 @@
Op o,
DataType ty,
builtin::AddressSpace as,
- const Symbol& buf)
- : Base(pid, nid), op(o), type(ty), address_space(as), buffer(buf) {}
+ const ast::IdentifierExpression* buf)
+ : Base(pid, nid, utils::Vector{buf}), op(o), type(ty), address_space(as) {}
DecomposeMemoryAccess::Intrinsic::~Intrinsic() = default;
std::string DecomposeMemoryAccess::Intrinsic::InternalName() const {
utils::StringStream ss;
@@ -794,7 +794,7 @@
const DecomposeMemoryAccess::Intrinsic* DecomposeMemoryAccess::Intrinsic::Clone(
CloneContext* ctx) const {
- auto buf = ctx->Clone(buffer);
+ auto buf = ctx->Clone(Buffer());
return ctx->dst->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
ctx->dst->ID(), ctx->dst->AllocateNodeID(), op, type, address_space, buf);
}
@@ -803,6 +803,10 @@
return op != Op::kLoad && op != Op::kStore;
}
+const ast::IdentifierExpression* DecomposeMemoryAccess::Intrinsic::Buffer() const {
+ return dependencies[0];
+}
+
DecomposeMemoryAccess::DecomposeMemoryAccess() = default;
DecomposeMemoryAccess::~DecomposeMemoryAccess() = default;
diff --git a/src/tint/transform/decompose_memory_access.h b/src/tint/transform/decompose_memory_access.h
index e34cd63..f85ad6d 100644
--- a/src/tint/transform/decompose_memory_access.h
+++ b/src/tint/transform/decompose_memory_access.h
@@ -80,13 +80,13 @@
/// @param o the op of the intrinsic
/// @param type the data type of the intrinsic
/// @param address_space the address space of the buffer
- /// @param buffer the storage or uniform buffer name
+ /// @param buffer the storage or uniform buffer identifier
Intrinsic(ProgramID pid,
ast::NodeID nid,
Op o,
DataType type,
builtin::AddressSpace address_space,
- const Symbol& buffer);
+ const ast::IdentifierExpression* buffer);
/// Destructor
~Intrinsic() override;
@@ -102,6 +102,9 @@
/// @return true if op is atomic
bool IsAtomic() const;
+ /// @return the buffer that this intrinsic operates on
+ const ast::IdentifierExpression* Buffer() const;
+
/// The op of the intrinsic
const Op op;
@@ -110,9 +113,6 @@
/// The address space of the buffer this intrinsic operates on
const builtin::AddressSpace address_space;
-
- /// The buffer name
- const Symbol buffer;
};
/// Constructor
diff --git a/src/tint/transform/spirv_atomic.cc b/src/tint/transform/spirv_atomic.cc
index b3924cc..043861b 100644
--- a/src/tint/transform/spirv_atomic.cc
+++ b/src/tint/transform/spirv_atomic.cc
@@ -294,7 +294,7 @@
SpirvAtomic::~SpirvAtomic() = default;
SpirvAtomic::Stub::Stub(ProgramID pid, ast::NodeID nid, sem::BuiltinType b)
- : Base(pid, nid), builtin(b) {}
+ : Base(pid, nid, utils::Empty), builtin(b) {}
SpirvAtomic::Stub::~Stub() = default;
std::string SpirvAtomic::Stub::InternalName() const {
return "@internal(spirv-atomic " + std::string(sem::str(builtin)) + ")";
diff --git a/src/tint/utils/unique_vector.h b/src/tint/utils/unique_vector.h
index 6cb8f88..eca0d4c 100644
--- a/src/tint/utils/unique_vector.h
+++ b/src/tint/utils/unique_vector.h
@@ -30,6 +30,9 @@
/// Attempting to add a duplicate is a no-op.
template <typename T, size_t N, typename HASH = std::hash<T>, typename EQUAL = std::equal_to<T>>
struct UniqueVector {
+ /// STL-friendly alias to T. Used by gmock.
+ using value_type = T;
+
/// Constructor
UniqueVector() = default;
diff --git a/src/tint/writer/append_vector.cc b/src/tint/writer/append_vector.cc
index c31667d..98d9798 100644
--- a/src/tint/writer/append_vector.cc
+++ b/src/tint/writer/append_vector.cc
@@ -156,13 +156,12 @@
}));
auto* ctor_target = b->create<sem::ValueConstructor>(
packed_sem_ty,
- utils::Transform(
- packed,
- [&](const tint::sem::ValueExpression* arg, size_t i) -> const sem::Parameter* {
- return b->create<sem::Parameter>(
- nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
- builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
- }),
+ utils::Transform(packed,
+ [&](const tint::sem::ValueExpression* arg, size_t i) {
+ return b->create<sem::Parameter>(
+ nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
+ builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ }),
sem::EvaluationStage::kRuntime);
auto* ctor_sem = b->create<sem::Call>(ctor_ast, ctor_target, sem::EvaluationStage::kRuntime,
std::move(packed), statement,
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 823e6bd..53f8085 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1125,7 +1125,7 @@
utils::StringStream& out,
const ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto* const offset = expr->args[0];
// offset in bytes
@@ -1413,7 +1413,7 @@
utils::StringStream& out,
const ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto* const offset = expr->args[0];
auto* const value = expr->args[1];
@@ -1581,7 +1581,7 @@
const auto name = builder_.Symbols().NameFor(func->name->symbol);
auto& buf = *current_buffer_;
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto rmw = [&](const char* hlsl) -> bool {
{
diff --git a/test/tint/bug/tint/1860.wgsl b/test/tint/bug/tint/1860.wgsl
new file mode 100644
index 0000000..7fdd68f
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl
@@ -0,0 +1,10 @@
+@vertex
+fn main() -> @builtin(position) vec4<f32> {
+ return vec4(declared_after_usage.f);
+}
+
+struct DeclaredAfterUsage {
+ f : f32,
+}
+
+@group(0) @binding(0) var <uniform> declared_after_usage : DeclaredAfterUsage;
diff --git a/test/tint/bug/tint/1860.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1860.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..f89c071
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+cbuffer cbuffer_declared_after_usage : register(b0, space0) {
+ uint4 declared_after_usage[1];
+};
+
+float4 main_inner() {
+ return float4((asfloat(declared_after_usage[0].x)).xxxx);
+}
+
+tint_symbol main() {
+ const float4 inner_result = main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
diff --git a/test/tint/bug/tint/1860.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1860.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..f89c071
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+cbuffer cbuffer_declared_after_usage : register(b0, space0) {
+ uint4 declared_after_usage[1];
+};
+
+float4 main_inner() {
+ return float4((asfloat(declared_after_usage[0].x)).xxxx);
+}
+
+tint_symbol main() {
+ const float4 inner_result = main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
diff --git a/test/tint/bug/tint/1860.wgsl.expected.glsl b/test/tint/bug/tint/1860.wgsl.expected.glsl
new file mode 100644
index 0000000..b7b713a
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.glsl
@@ -0,0 +1,25 @@
+#version 310 es
+
+struct DeclaredAfterUsage {
+ float f;
+ uint pad;
+ uint pad_1;
+ uint pad_2;
+};
+
+layout(binding = 0, std140) uniform declared_after_usage_block_ubo {
+ DeclaredAfterUsage inner;
+} declared_after_usage;
+
+vec4 tint_symbol() {
+ return vec4(declared_after_usage.inner.f);
+}
+
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = tint_symbol();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
diff --git a/test/tint/bug/tint/1860.wgsl.expected.msl b/test/tint/bug/tint/1860.wgsl.expected.msl
new file mode 100644
index 0000000..98df1c9
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.msl
@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ float4 value [[position]];
+};
+
+struct DeclaredAfterUsage {
+ /* 0x0000 */ float f;
+};
+
+float4 tint_symbol_inner(const constant DeclaredAfterUsage* const tint_symbol_2) {
+ return float4((*(tint_symbol_2)).f);
+}
+
+vertex tint_symbol_1 tint_symbol(const constant DeclaredAfterUsage* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = tint_symbol_inner(tint_symbol_3);
+ tint_symbol_1 wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
diff --git a/test/tint/bug/tint/1860.wgsl.expected.spvasm b/test/tint/bug/tint/1860.wgsl.expected.spvasm
new file mode 100644
index 0000000..6577651
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.spvasm
@@ -0,0 +1,58 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %main "main" %value %vertex_point_size
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %declared_after_usage_block "declared_after_usage_block"
+ OpMemberName %declared_after_usage_block 0 "inner"
+ OpName %DeclaredAfterUsage "DeclaredAfterUsage"
+ OpMemberName %DeclaredAfterUsage 0 "f"
+ OpName %declared_after_usage "declared_after_usage"
+ OpName %main_inner "main_inner"
+ OpName %main "main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %declared_after_usage_block Block
+ OpMemberDecorate %declared_after_usage_block 0 Offset 0
+ OpMemberDecorate %DeclaredAfterUsage 0 Offset 0
+ OpDecorate %declared_after_usage NonWritable
+ OpDecorate %declared_after_usage DescriptorSet 0
+ OpDecorate %declared_after_usage Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+%DeclaredAfterUsage = OpTypeStruct %float
+%declared_after_usage_block = OpTypeStruct %DeclaredAfterUsage
+%_ptr_Uniform_declared_after_usage_block = OpTypePointer Uniform %declared_after_usage_block
+%declared_after_usage = OpVariable %_ptr_Uniform_declared_after_usage_block Uniform
+ %13 = OpTypeFunction %v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %void = OpTypeVoid
+ %22 = OpTypeFunction %void
+ %float_1 = OpConstant %float 1
+ %main_inner = OpFunction %v4float None %13
+ %15 = OpLabel
+ %19 = OpAccessChain %_ptr_Uniform_float %declared_after_usage %uint_0 %uint_0
+ %20 = OpLoad %float %19
+ %21 = OpCompositeConstruct %v4float %20 %20 %20 %20
+ OpReturnValue %21
+ OpFunctionEnd
+ %main = OpFunction %void None %22
+ %25 = OpLabel
+ %26 = OpFunctionCall %v4float %main_inner
+ OpStore %value %26
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/1860.wgsl.expected.wgsl b/test/tint/bug/tint/1860.wgsl.expected.wgsl
new file mode 100644
index 0000000..562b411
--- /dev/null
+++ b/test/tint/bug/tint/1860.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+@vertex
+fn main() -> @builtin(position) vec4<f32> {
+ return vec4(declared_after_usage.f);
+}
+
+struct DeclaredAfterUsage {
+ f : f32,
+}
+
+@group(0) @binding(0) var<uniform> declared_after_usage : DeclaredAfterUsage;
