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>

src/tint/ast/disable_validation_attribute.cc

@@ -23,7 +23,7 @@ namespace tint::ast {
 DisableValidationAttribute::DisableValidationAttribute(ProgramID pid,
                                                        NodeID nid,
                                                        DisabledValidation val)
-    : Base(pid, nid), validation(val) {}
+    : Base(pid, nid, utils::Empty), validation(val) {}
 
 DisableValidationAttribute::~DisableValidationAttribute() = default;
 

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;
 

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 @@ class InternalAttribute : public Castable<InternalAttribute, Attribute> {
     /// 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 @@ class InternalAttribute : public Castable<InternalAttribute, Attribute> {
 
     /// @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

src/tint/resolver/attribute_validation_test.cc

@@ -1977,4 +1977,40 @@ TEST_F(MustUseAttributeTest, UsedOnFnWithNoReturnValue) {
 }  // 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);

src/tint/resolver/dependency_graph.cc

@@ -414,12 +414,18 @@ class DependencyScanner {
                 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>()) {

src/tint/resolver/intrinsic_table.cc

@@ -1495,7 +1495,7 @@ IntrinsicTable::CtorOrConv Impl::Lookup(CtorConvIntrinsic type,
 
     // 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>(

src/tint/resolver/resolver.cc

@@ -856,9 +856,9 @@ sem::Statement* Resolver::ConstAssert(const ast::ConstAssert* assertion) {
 sem::Function* Resolver::Function(const ast::Function* decl) {
     Mark(decl->name);
 
-    uint32_t parameter_index = 0;
+    auto* func = builder_->create<sem::Function>(decl);
-    utils::Hashmap<Symbol, Source, 8> parameter_names;
+    builder_->Sem().Add(decl, func);
-    utils::Vector<sem::Parameter*, 8> parameters;
+    TINT_SCOPED_ASSIGNMENT(current_function_, func);
 
     validator_.DiagnosticFilters().Push();
     TINT_DEFER(validator_.DiagnosticFilters().Pop());
@@ -872,6 +872,8 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
     }
 
     // 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 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
             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 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
     } 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 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
             if (!value) {
                 return nullptr;
             }
-            return_location = value.Get();
+            func->SetReturnLocation(value.Get());
         }
     }
 
@@ -963,12 +965,7 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
         }
     }
 
-    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 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
                 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 @@ bool Resolver::Attribute(const ast::Attribute* attr) {
         [&](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 @@ bool Resolver::InterpolateAttribute(const ast::InterpolateAttribute* attr) {
     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);
 

src/tint/resolver/resolver.h

@@ -321,6 +321,10 @@ class Resolver {
     /// @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);

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 @@ Builtin::Builtin(BuiltinType type,
                  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) {}

src/tint/sem/call_target.cc

@@ -23,17 +23,25 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::CallTarget);
 
 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)) {}

src/tint/sem/call_target.h

@@ -27,6 +27,9 @@ namespace tint::sem {
 
 /// 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
@@ -39,9 +42,9 @@ struct CallTargetSignature {
     ~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
@@ -66,6 +69,12 @@ struct CallTargetSignature {
 /// conversions.
 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
@@ -73,7 +82,7 @@ class CallTarget : public Castable<CallTarget, Node> {
     /// @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 @@ class CallTarget : public Castable<CallTarget, Node> {
     /// 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; }
 

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,
+Function::Function(const ast::Function* declaration)
-                                            const tint::sem::CallTarget* owner) {
+    : Base(EvaluationStage::kRuntime,
-    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,
            ast::HasAttribute<ast::MustUseAttribute>(declaration->attributes)),
       declaration_(declaration),
-      workgroup_size_{1, 1, 1},
+      workgroup_size_{1, 1, 1} {}
-      return_location_(return_location) {}
 
 Function::~Function() = default;
 

src/tint/sem/function.h

@@ -54,17 +54,15 @@ class Function final : public Castable<Function, CallTarget> {
 
     /// Constructor
     /// @param declaration the ast::Function
-    /// @param return_type the return type of the function
+    explicit Function(const ast::Function* declaration);
-    /// @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);
 
     /// 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_; }
 

src/tint/sem/value_constructor.cc

@@ -21,7 +21,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::ValueConstructor);
 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) {}
 

src/tint/sem/value_constructor.h

@@ -28,7 +28,7 @@ class ValueConstructor final : public Castable<ValueConstructor, CallTarget> {
     /// @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

src/tint/sem/value_conversion.cc

@@ -19,9 +19,9 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::ValueConversion);
 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;
 

src/tint/sem/value_conversion.h

@@ -26,7 +26,7 @@ class ValueConversion final : public Castable<ValueConversion, CallTarget> {
     /// @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;

src/tint/transform/add_block_attribute.cc

@@ -102,7 +102,7 @@ Transform::ApplyResult AddBlockAttribute::Apply(const Program* src,
 }
 
 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";

src/tint/transform/calculate_array_length.cc

@@ -71,7 +71,7 @@ struct ArrayUsage {
 }  // 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";

src/tint/transform/decompose_memory_access.cc

@@ -228,7 +228,7 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicLoadFor(ProgramBuilder* builder,
     }
     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 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicStoreFor(ProgramBuilder* builder,
     }
     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 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicAtomicFor(ProgramBuilder* builder,
     }
     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 @@ DecomposeMemoryAccess::Intrinsic::Intrinsic(ProgramID pid,
                                             Op o,
                                             DataType ty,
                                             builtin::AddressSpace as,
-                                            const Symbol& buf)
+                                            const ast::IdentifierExpression* buf)
-    : Base(pid, nid), op(o), type(ty), address_space(as), buffer(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 @@ std::string DecomposeMemoryAccess::Intrinsic::InternalName() const {
 
 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 @@ bool DecomposeMemoryAccess::Intrinsic::IsAtomic() const {
     return op != Op::kLoad && op != Op::kStore;
 }
 
+const ast::IdentifierExpression* DecomposeMemoryAccess::Intrinsic::Buffer() const {
+    return dependencies[0];
+}
+
 DecomposeMemoryAccess::DecomposeMemoryAccess() = default;
 DecomposeMemoryAccess::~DecomposeMemoryAccess() = default;
 

src/tint/transform/decompose_memory_access.h

@@ -80,13 +80,13 @@ class DecomposeMemoryAccess final : public Castable<DecomposeMemoryAccess, Trans
         /// @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 @@ class DecomposeMemoryAccess final : public Castable<DecomposeMemoryAccess, Trans
         /// @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 @@ class DecomposeMemoryAccess final : public Castable<DecomposeMemoryAccess, Trans
 
         /// The address space of the buffer this intrinsic operates on
         const builtin::AddressSpace address_space;
-
-        /// The buffer name
-        const Symbol buffer;
     };
 
     /// Constructor

src/tint/transform/spirv_atomic.cc

@@ -294,7 +294,7 @@ SpirvAtomic::SpirvAtomic() = default;
 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)) + ")";

src/tint/utils/unique_vector.h

@@ -30,6 +30,9 @@ namespace tint::utils {
 /// 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;
 

src/tint/writer/append_vector.cc

@@ -156,13 +156,12 @@ const sem::Call* AppendVector(ProgramBuilder* b,
                 }));
     auto* ctor_target = b->create<sem::ValueConstructor>(
         packed_sem_ty,
-        utils::Transform(
+        utils::Transform(packed,
-            packed,
+                         [&](const tint::sem::ValueExpression* arg, size_t i) {
-            [&](const tint::sem::ValueExpression* arg, size_t i) -> const sem::Parameter* {
+                             return b->create<sem::Parameter>(
-                return b->create<sem::Parameter>(
+                                 nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
-                    nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
+                                 builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
-                    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,

src/tint/writer/hlsl/generator_impl.cc

@@ -1125,7 +1125,7 @@ bool GeneratorImpl::EmitUniformBufferAccess(
     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 @@ bool GeneratorImpl::EmitStorageBufferAccess(
     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 @@ bool GeneratorImpl::EmitStorageAtomicIntrinsic(
     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 {
         {

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;

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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;