Import Tint changes from Dawn

Changes:
  - c4e076ffe6365b8e5254488362340cfbb355efc6 Rename attribute values. by dan sinclair <dsinclair@chromium.org>
  - f9eeed61067764fb78f63a106dbda4e6a47e5787 Convert `@location` to store expression internally. by dan sinclair <dsinclair@chromium.org>
  - 145337f309abad41995dbd6740976c72b9d2968f Use SubstituteOverride transform to implement overrides by shrekshao <shrekshao@google.com>
  - 23cf74c30ea598dd18682eb041217f0e34bd46ad Allow sem::GlobalVariable to hold a location. by dan sinclair <dsinclair@chromium.org>
  - 29aa613dcfb453270335f146031bbc084542923f tint: const eval of comparison operations by Antonio Maiorano <amaiorano@google.com>
  - eeda18d55ea2f998b5eaff583b4e1922891f8013 tint::transform::SingleEntryPoint: Preserve global 'const's by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: c4e076ffe6365b8e5254488362340cfbb355efc6
Change-Id: Ide1598f6a2a04b4d37de6dcf972e3a11efa22437
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/101660
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/ast/binding_attribute.cc b/src/tint/ast/binding_attribute.cc
index 405fe98..38f1d0f 100644
--- a/src/tint/ast/binding_attribute.cc
+++ b/src/tint/ast/binding_attribute.cc
@@ -25,8 +25,8 @@
 BindingAttribute::BindingAttribute(ProgramID pid,
                                    NodeID nid,
                                    const Source& src,
-                                   const ast::Expression* val)
-    : Base(pid, nid, src), value(val) {}
+                                   const ast::Expression* exp)
+    : Base(pid, nid, src), expr(exp) {}
 
 BindingAttribute::~BindingAttribute() = default;
 
@@ -37,8 +37,8 @@
 const BindingAttribute* BindingAttribute::Clone(CloneContext* ctx) const {
     // Clone arguments outside of create() call to have deterministic ordering
     auto src = ctx->Clone(source);
-    auto* value_ = ctx->Clone(value);
-    return ctx->dst->create<BindingAttribute>(src, value_);
+    auto* expr_ = ctx->Clone(expr);
+    return ctx->dst->create<BindingAttribute>(src, expr_);
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/binding_attribute.h b/src/tint/ast/binding_attribute.h
index 39ca2b6..7bb7add 100644
--- a/src/tint/ast/binding_attribute.h
+++ b/src/tint/ast/binding_attribute.h
@@ -29,8 +29,8 @@
     /// @param pid the identifier of the program that owns this node
     /// @param nid the unique node identifier
     /// @param src the source of this node
-    /// @param value the binding value expression
-    BindingAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* value);
+    /// @param expr the binding expression
+    BindingAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* expr);
     ~BindingAttribute() override;
 
     /// @returns the WGSL name for the attribute
@@ -42,8 +42,8 @@
     /// @return the newly cloned node
     const BindingAttribute* Clone(CloneContext* ctx) const override;
 
-    /// the binding value expression
-    const ast::Expression* const value;
+    /// the binding expression
+    const ast::Expression* const expr;
 };
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/binding_attribute_test.cc b/src/tint/ast/binding_attribute_test.cc
index ec8538c..941bd68 100644
--- a/src/tint/ast/binding_attribute_test.cc
+++ b/src/tint/ast/binding_attribute_test.cc
@@ -22,7 +22,7 @@
 
 TEST_F(BindingAttributeTest, Creation) {
     auto* d = Binding(2_a);
-    EXPECT_TRUE(d->value->Is<IntLiteralExpression>());
+    EXPECT_TRUE(d->expr->Is<IntLiteralExpression>());
 }
 
 }  // namespace
diff --git a/src/tint/ast/group_attribute.cc b/src/tint/ast/group_attribute.cc
index 9f44003..ff60bd5 100644
--- a/src/tint/ast/group_attribute.cc
+++ b/src/tint/ast/group_attribute.cc
@@ -25,8 +25,8 @@
 GroupAttribute::GroupAttribute(ProgramID pid,
                                NodeID nid,
                                const Source& src,
-                               const ast::Expression* val)
-    : Base(pid, nid, src), value(val) {}
+                               const ast::Expression* exp)
+    : Base(pid, nid, src), expr(exp) {}
 
 GroupAttribute::~GroupAttribute() = default;
 
@@ -37,8 +37,8 @@
 const GroupAttribute* GroupAttribute::Clone(CloneContext* ctx) const {
     // Clone arguments outside of create() call to have deterministic ordering
     auto src = ctx->Clone(source);
-    auto* value_ = ctx->Clone(value);
-    return ctx->dst->create<GroupAttribute>(src, value_);
+    auto* expr_ = ctx->Clone(expr);
+    return ctx->dst->create<GroupAttribute>(src, expr_);
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/group_attribute.h b/src/tint/ast/group_attribute.h
index 5110489..552a69f 100644
--- a/src/tint/ast/group_attribute.h
+++ b/src/tint/ast/group_attribute.h
@@ -29,8 +29,8 @@
     /// @param pid the identifier of the program that owns this node
     /// @param nid the unique node identifier
     /// @param src the source of this node
-    /// @param value the group value expression
-    GroupAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* value);
+    /// @param expr the group expression
+    GroupAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* expr);
     ~GroupAttribute() override;
 
     /// @returns the WGSL name for the attribute
@@ -42,8 +42,8 @@
     /// @return the newly cloned node
     const GroupAttribute* Clone(CloneContext* ctx) const override;
 
-    /// The group value expression
-    const ast::Expression* const value;
+    /// The group expression
+    const ast::Expression* const expr;
 };
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/group_attribute_test.cc b/src/tint/ast/group_attribute_test.cc
index 616ea37..c8e2130 100644
--- a/src/tint/ast/group_attribute_test.cc
+++ b/src/tint/ast/group_attribute_test.cc
@@ -22,7 +22,7 @@
 
 TEST_F(GroupAttributeTest, Creation) {
     auto* d = Group(2_a);
-    EXPECT_TRUE(d->value->Is<IntLiteralExpression>());
+    EXPECT_TRUE(d->expr->Is<IntLiteralExpression>());
 }
 
 }  // namespace
diff --git a/src/tint/ast/id_attribute.cc b/src/tint/ast/id_attribute.cc
index 9c1d1ae..0515f01 100644
--- a/src/tint/ast/id_attribute.cc
+++ b/src/tint/ast/id_attribute.cc
@@ -22,8 +22,8 @@
 
 namespace tint::ast {
 
-IdAttribute::IdAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* val)
-    : Base(pid, nid, src), value(val) {}
+IdAttribute::IdAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* exp)
+    : Base(pid, nid, src), expr(exp) {}
 
 IdAttribute::~IdAttribute() = default;
 
@@ -34,8 +34,8 @@
 const IdAttribute* IdAttribute::Clone(CloneContext* ctx) const {
     // Clone arguments outside of create() call to have deterministic ordering
     auto src = ctx->Clone(source);
-    auto* value_ = ctx->Clone(value);
-    return ctx->dst->create<IdAttribute>(src, value_);
+    auto* expr_ = ctx->Clone(expr);
+    return ctx->dst->create<IdAttribute>(src, expr_);
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/id_attribute.h b/src/tint/ast/id_attribute.h
index f707bde..f683080 100644
--- a/src/tint/ast/id_attribute.h
+++ b/src/tint/ast/id_attribute.h
@@ -29,8 +29,8 @@
     /// @param pid the identifier of the program that owns this node
     /// @param nid the unique node identifier
     /// @param src the source of this node
-    /// @param val the numeric id value expression
-    IdAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* val);
+    /// @param expr the numeric id expression
+    IdAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* expr);
     ~IdAttribute() override;
 
     /// @returns the WGSL name for the attribute
@@ -42,8 +42,8 @@
     /// @return the newly cloned node
     const IdAttribute* Clone(CloneContext* ctx) const override;
 
-    /// The id value expression
-    const ast::Expression* const value;
+    /// The id expression
+    const ast::Expression* const expr;
 };
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/id_attribute_test.cc b/src/tint/ast/id_attribute_test.cc
index 84605b1..eeef23b 100644
--- a/src/tint/ast/id_attribute_test.cc
+++ b/src/tint/ast/id_attribute_test.cc
@@ -24,7 +24,7 @@
 
 TEST_F(IdAttributeTest, Creation) {
     auto* d = Id(12_a);
-    EXPECT_TRUE(d->value->Is<ast::IntLiteralExpression>());
+    EXPECT_TRUE(d->expr->Is<ast::IntLiteralExpression>());
 }
 
 }  // namespace
diff --git a/src/tint/ast/location_attribute.cc b/src/tint/ast/location_attribute.cc
index 2ea2d5d..4f34144 100644
--- a/src/tint/ast/location_attribute.cc
+++ b/src/tint/ast/location_attribute.cc
@@ -22,8 +22,11 @@
 
 namespace tint::ast {
 
-LocationAttribute::LocationAttribute(ProgramID pid, NodeID nid, const Source& src, uint32_t val)
-    : Base(pid, nid, src), value(val) {}
+LocationAttribute::LocationAttribute(ProgramID pid,
+                                     NodeID nid,
+                                     const Source& src,
+                                     const ast::Expression* exp)
+    : Base(pid, nid, src), expr(exp) {}
 
 LocationAttribute::~LocationAttribute() = default;
 
@@ -34,7 +37,8 @@
 const LocationAttribute* LocationAttribute::Clone(CloneContext* ctx) const {
     // Clone arguments outside of create() call to have deterministic ordering
     auto src = ctx->Clone(source);
-    return ctx->dst->create<LocationAttribute>(src, value);
+    auto expr_ = ctx->Clone(expr);
+    return ctx->dst->create<LocationAttribute>(src, expr_);
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/location_attribute.h b/src/tint/ast/location_attribute.h
index 97c6fea..43d5edf 100644
--- a/src/tint/ast/location_attribute.h
+++ b/src/tint/ast/location_attribute.h
@@ -18,6 +18,7 @@
 #include <string>
 
 #include "src/tint/ast/attribute.h"
+#include "src/tint/ast/expression.h"
 
 namespace tint::ast {
 
@@ -28,8 +29,8 @@
     /// @param pid the identifier of the program that owns this node
     /// @param nid the unique node identifier
     /// @param src the source of this node
-    /// @param value the location value
-    LocationAttribute(ProgramID pid, NodeID nid, const Source& src, uint32_t value);
+    /// @param expr the location expression
+    LocationAttribute(ProgramID pid, NodeID nid, const Source& src, const ast::Expression* expr);
     ~LocationAttribute() override;
 
     /// @returns the WGSL name for the attribute
@@ -41,8 +42,8 @@
     /// @return the newly cloned node
     const LocationAttribute* Clone(CloneContext* ctx) const override;
 
-    /// The location value
-    const uint32_t value;
+    /// The location expression
+    const ast::Expression* const expr;
 };
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/location_attribute_test.cc b/src/tint/ast/location_attribute_test.cc
index e0bcb39..681ea43 100644
--- a/src/tint/ast/location_attribute_test.cc
+++ b/src/tint/ast/location_attribute_test.cc
@@ -17,11 +17,12 @@
 namespace tint::ast {
 namespace {
 
+using namespace tint::number_suffixes;  // NOLINT
 using LocationAttributeTest = TestHelper;
 
 TEST_F(LocationAttributeTest, Creation) {
-    auto* d = create<LocationAttribute>(2u);
-    EXPECT_EQ(2u, d->value);
+    auto* d = Location(2_a);
+    EXPECT_TRUE(d->expr->Is<IntLiteralExpression>());
 }
 
 }  // namespace
diff --git a/src/tint/ast/struct_member_align_attribute.cc b/src/tint/ast/struct_member_align_attribute.cc
index e188e7b..2a21362 100644
--- a/src/tint/ast/struct_member_align_attribute.cc
+++ b/src/tint/ast/struct_member_align_attribute.cc
@@ -27,7 +27,7 @@
                                                        NodeID nid,
                                                        const Source& src,
                                                        const ast::Expression* a)
-    : Base(pid, nid, src), align(a) {}
+    : Base(pid, nid, src), expr(a) {}
 
 StructMemberAlignAttribute::~StructMemberAlignAttribute() = default;
 
@@ -38,8 +38,8 @@
 const StructMemberAlignAttribute* StructMemberAlignAttribute::Clone(CloneContext* ctx) const {
     // Clone arguments outside of create() call to have deterministic ordering
     auto src = ctx->Clone(source);
-    auto* align_ = ctx->Clone(align);
-    return ctx->dst->create<StructMemberAlignAttribute>(src, align_);
+    auto* expr_ = ctx->Clone(expr);
+    return ctx->dst->create<StructMemberAlignAttribute>(src, expr_);
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/struct_member_align_attribute.h b/src/tint/ast/struct_member_align_attribute.h
index 2043b01..6da1894 100644
--- a/src/tint/ast/struct_member_align_attribute.h
+++ b/src/tint/ast/struct_member_align_attribute.h
@@ -30,7 +30,7 @@
     /// @param pid the identifier of the program that owns this node
     /// @param nid the unique node identifier
     /// @param src the source of this node
-    /// @param align the align value expression
+    /// @param align the align expression
     StructMemberAlignAttribute(ProgramID pid,
                                NodeID nid,
                                const Source& src,
@@ -46,8 +46,8 @@
     /// @return the newly cloned node
     const StructMemberAlignAttribute* Clone(CloneContext* ctx) const override;
 
-    /// The align value expression
-    const ast::Expression* const align;
+    /// The align expression
+    const ast::Expression* const expr;
 };
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/struct_member_align_attribute_test.cc b/src/tint/ast/struct_member_align_attribute_test.cc
index f52d32b..3915a6e 100644
--- a/src/tint/ast/struct_member_align_attribute_test.cc
+++ b/src/tint/ast/struct_member_align_attribute_test.cc
@@ -24,8 +24,8 @@
 TEST_F(StructMemberAlignAttributeTest, Creation) {
     auto* val = Expr("ident");
     auto* d = create<StructMemberAlignAttribute>(val);
-    EXPECT_EQ(val, d->align);
-    EXPECT_TRUE(d->align->Is<IdentifierExpression>());
+    EXPECT_EQ(val, d->expr);
+    EXPECT_TRUE(d->expr->Is<IdentifierExpression>());
 }
 
 }  // namespace
diff --git a/src/tint/ast/variable_test.cc b/src/tint/ast/variable_test.cc
index 40dd68d..2fed042 100644
--- a/src/tint/ast/variable_test.cc
+++ b/src/tint/ast/variable_test.cc
@@ -92,7 +92,7 @@
 }
 
 TEST_F(VariableTest, WithAttributes) {
-    auto* var = Var("my_var", ty.i32(), StorageClass::kFunction, Location(1u),
+    auto* var = Var("my_var", ty.i32(), StorageClass::kFunction, Location(1_u),
                     Builtin(BuiltinValue::kPosition), Id(1200_u));
 
     auto& attributes = var->attributes;
@@ -102,7 +102,8 @@
 
     auto* location = ast::GetAttribute<ast::LocationAttribute>(attributes);
     ASSERT_NE(nullptr, location);
-    EXPECT_EQ(1u, location->value);
+    ASSERT_NE(nullptr, location->expr);
+    EXPECT_TRUE(location->expr->Is<ast::IntLiteralExpression>());
 }
 
 TEST_F(VariableTest, HasBindingPoint_BothProvided) {
diff --git a/src/tint/inspector/inspector.cc b/src/tint/inspector/inspector.cc
index 2f8d09a..087e786 100644
--- a/src/tint/inspector/inspector.cc
+++ b/src/tint/inspector/inspector.cc
@@ -133,6 +133,110 @@
 
 Inspector::~Inspector() = default;
 
+EntryPoint Inspector::GetEntryPoint(const tint::ast::Function* func) {
+    EntryPoint entry_point;
+    TINT_ASSERT(Inspector, func != nullptr);
+    TINT_ASSERT(Inspector, func->IsEntryPoint());
+
+    auto* sem = program_->Sem().Get(func);
+
+    entry_point.name = program_->Symbols().NameFor(func->symbol);
+    entry_point.remapped_name = program_->Symbols().NameFor(func->symbol);
+
+    switch (func->PipelineStage()) {
+        case ast::PipelineStage::kCompute: {
+            entry_point.stage = PipelineStage::kCompute;
+
+            auto wgsize = sem->WorkgroupSize();
+            if (!wgsize[0].overridable_const && !wgsize[1].overridable_const &&
+                !wgsize[2].overridable_const) {
+                entry_point.workgroup_size = {wgsize[0].value, wgsize[1].value, wgsize[2].value};
+            }
+            break;
+        }
+        case ast::PipelineStage::kFragment: {
+            entry_point.stage = PipelineStage::kFragment;
+            break;
+        }
+        case ast::PipelineStage::kVertex: {
+            entry_point.stage = PipelineStage::kVertex;
+            break;
+        }
+        default: {
+            TINT_UNREACHABLE(Inspector, diagnostics_)
+                << "invalid pipeline stage for entry point '" << entry_point.name << "'";
+            break;
+        }
+    }
+
+    for (auto* param : sem->Parameters()) {
+        AddEntryPointInOutVariables(program_->Symbols().NameFor(param->Declaration()->symbol),
+                                    param->Type(), param->Declaration()->attributes,
+                                    param->Location(), entry_point.input_variables);
+
+        entry_point.input_position_used |= ContainsBuiltin(
+            ast::BuiltinValue::kPosition, param->Type(), param->Declaration()->attributes);
+        entry_point.front_facing_used |= ContainsBuiltin(
+            ast::BuiltinValue::kFrontFacing, param->Type(), param->Declaration()->attributes);
+        entry_point.sample_index_used |= ContainsBuiltin(
+            ast::BuiltinValue::kSampleIndex, param->Type(), param->Declaration()->attributes);
+        entry_point.input_sample_mask_used |= ContainsBuiltin(
+            ast::BuiltinValue::kSampleMask, param->Type(), param->Declaration()->attributes);
+        entry_point.num_workgroups_used |= ContainsBuiltin(
+            ast::BuiltinValue::kNumWorkgroups, param->Type(), param->Declaration()->attributes);
+    }
+
+    if (!sem->ReturnType()->Is<sem::Void>()) {
+        AddEntryPointInOutVariables("<retval>", sem->ReturnType(), func->return_type_attributes,
+                                    sem->ReturnLocation(), entry_point.output_variables);
+
+        entry_point.output_sample_mask_used = ContainsBuiltin(
+            ast::BuiltinValue::kSampleMask, sem->ReturnType(), func->return_type_attributes);
+    }
+
+    for (auto* var : sem->TransitivelyReferencedGlobals()) {
+        auto* decl = var->Declaration();
+
+        auto name = program_->Symbols().NameFor(decl->symbol);
+
+        auto* global = var->As<sem::GlobalVariable>();
+        if (global && global->Declaration()->Is<ast::Override>()) {
+            Override override;
+            override.name = name;
+            override.id = global->OverrideId();
+            auto* type = var->Type();
+            TINT_ASSERT(Inspector, type->is_scalar());
+            if (type->is_bool_scalar_or_vector()) {
+                override.type = Override::Type::kBool;
+            } else if (type->is_float_scalar()) {
+                override.type = Override::Type::kFloat32;
+            } else if (type->is_signed_integer_scalar()) {
+                override.type = Override::Type::kInt32;
+            } else if (type->is_unsigned_integer_scalar()) {
+                override.type = Override::Type::kUint32;
+            } else {
+                TINT_UNREACHABLE(Inspector, diagnostics_);
+            }
+
+            override.is_initialized = global->Declaration()->constructor;
+            override.is_id_specified =
+                ast::HasAttribute<ast::IdAttribute>(global->Declaration()->attributes);
+
+            entry_point.overrides.push_back(override);
+        }
+    }
+
+    return entry_point;
+}
+
+EntryPoint Inspector::GetEntryPoint(const std::string& entry_point_name) {
+    auto* func = FindEntryPointByName(entry_point_name);
+    if (!func) {
+        return EntryPoint();
+    }
+    return GetEntryPoint(func);
+}
+
 std::vector<EntryPoint> Inspector::GetEntryPoints() {
     std::vector<EntryPoint> result;
 
@@ -141,97 +245,7 @@
             continue;
         }
 
-        auto* sem = program_->Sem().Get(func);
-
-        EntryPoint entry_point;
-        entry_point.name = program_->Symbols().NameFor(func->symbol);
-        entry_point.remapped_name = program_->Symbols().NameFor(func->symbol);
-
-        switch (func->PipelineStage()) {
-            case ast::PipelineStage::kCompute: {
-                entry_point.stage = PipelineStage::kCompute;
-
-                auto wgsize = sem->WorkgroupSize();
-                if (!wgsize[0].overridable_const && !wgsize[1].overridable_const &&
-                    !wgsize[2].overridable_const) {
-                    entry_point.workgroup_size = {wgsize[0].value, wgsize[1].value,
-                                                  wgsize[2].value};
-                }
-                break;
-            }
-            case ast::PipelineStage::kFragment: {
-                entry_point.stage = PipelineStage::kFragment;
-                break;
-            }
-            case ast::PipelineStage::kVertex: {
-                entry_point.stage = PipelineStage::kVertex;
-                break;
-            }
-            default: {
-                TINT_UNREACHABLE(Inspector, diagnostics_)
-                    << "invalid pipeline stage for entry point '" << entry_point.name << "'";
-                break;
-            }
-        }
-
-        for (auto* param : sem->Parameters()) {
-            AddEntryPointInOutVariables(program_->Symbols().NameFor(param->Declaration()->symbol),
-                                        param->Type(), param->Declaration()->attributes,
-                                        entry_point.input_variables);
-
-            entry_point.input_position_used |= ContainsBuiltin(
-                ast::BuiltinValue::kPosition, param->Type(), param->Declaration()->attributes);
-            entry_point.front_facing_used |= ContainsBuiltin(
-                ast::BuiltinValue::kFrontFacing, param->Type(), param->Declaration()->attributes);
-            entry_point.sample_index_used |= ContainsBuiltin(
-                ast::BuiltinValue::kSampleIndex, param->Type(), param->Declaration()->attributes);
-            entry_point.input_sample_mask_used |= ContainsBuiltin(
-                ast::BuiltinValue::kSampleMask, param->Type(), param->Declaration()->attributes);
-            entry_point.num_workgroups_used |= ContainsBuiltin(
-                ast::BuiltinValue::kNumWorkgroups, param->Type(), param->Declaration()->attributes);
-        }
-
-        if (!sem->ReturnType()->Is<sem::Void>()) {
-            AddEntryPointInOutVariables("<retval>", sem->ReturnType(), func->return_type_attributes,
-                                        entry_point.output_variables);
-
-            entry_point.output_sample_mask_used = ContainsBuiltin(
-                ast::BuiltinValue::kSampleMask, sem->ReturnType(), func->return_type_attributes);
-        }
-
-        for (auto* var : sem->TransitivelyReferencedGlobals()) {
-            auto* decl = var->Declaration();
-
-            auto name = program_->Symbols().NameFor(decl->symbol);
-
-            auto* global = var->As<sem::GlobalVariable>();
-            if (global && global->Declaration()->Is<ast::Override>()) {
-                Override override;
-                override.name = name;
-                override.id = global->OverrideId();
-                auto* type = var->Type();
-                TINT_ASSERT(Inspector, type->is_scalar());
-                if (type->is_bool_scalar_or_vector()) {
-                    override.type = Override::Type::kBool;
-                } else if (type->is_float_scalar()) {
-                    override.type = Override::Type::kFloat32;
-                } else if (type->is_signed_integer_scalar()) {
-                    override.type = Override::Type::kInt32;
-                } else if (type->is_unsigned_integer_scalar()) {
-                    override.type = Override::Type::kUint32;
-                } else {
-                    TINT_UNREACHABLE(Inspector, diagnostics_);
-                }
-
-                override.is_initialized = global->Declaration()->constructor;
-                override.is_id_specified =
-                    ast::HasAttribute<ast::IdAttribute>(global->Declaration()->attributes);
-
-                entry_point.overrides.push_back(override);
-            }
-        }
-
-        result.push_back(std::move(entry_point));
+        result.push_back(GetEntryPoint(func));
     }
 
     return result;
@@ -609,6 +623,7 @@
 void Inspector::AddEntryPointInOutVariables(std::string name,
                                             const sem::Type* type,
                                             utils::VectorRef<const ast::Attribute*> attributes,
+                                            std::optional<uint32_t> location,
                                             std::vector<StageVariable>& variables) const {
     // Skip builtins.
     if (ast::HasAttribute<ast::BuiltinAttribute>(attributes)) {
@@ -622,7 +637,7 @@
         for (auto* member : struct_ty->Members()) {
             AddEntryPointInOutVariables(
                 name + "." + program_->Symbols().NameFor(member->Declaration()->symbol),
-                member->Type(), member->Declaration()->attributes, variables);
+                member->Type(), member->Declaration()->attributes, member->Location(), variables);
         }
         return;
     }
@@ -634,10 +649,9 @@
     std::tie(stage_variable.component_type, stage_variable.composition_type) =
         CalculateComponentAndComposition(type);
 
-    auto* location = ast::GetAttribute<ast::LocationAttribute>(attributes);
-    TINT_ASSERT(Inspector, location != nullptr);
+    TINT_ASSERT(Inspector, location.has_value());
     stage_variable.has_location_attribute = true;
-    stage_variable.location_attribute = location->value;
+    stage_variable.location_attribute = location.value();
 
     std::tie(stage_variable.interpolation_type, stage_variable.interpolation_sampling) =
         CalculateInterpolationData(type, attributes);
diff --git a/src/tint/inspector/inspector.h b/src/tint/inspector/inspector.h
index f3fe270..49e4bdf 100644
--- a/src/tint/inspector/inspector.h
+++ b/src/tint/inspector/inspector.h
@@ -55,6 +55,10 @@
     /// @returns vector of entry point information
     std::vector<EntryPoint> GetEntryPoints();
 
+    /// @param entry_point name of the entry point to get information about
+    /// @returns the entry point information
+    EntryPoint GetEntryPoint(const std::string& entry_point);
+
     /// @returns map of override identifier to initial value
     std::map<OverrideId, Scalar> GetOverrideDefaultValues();
 
@@ -168,10 +172,12 @@
     /// @param name the name of the variable being added
     /// @param type the type of the variable
     /// @param attributes the variable attributes
+    /// @param location the location value if provided
     /// @param variables the list to add the variables to
     void AddEntryPointInOutVariables(std::string name,
                                      const sem::Type* type,
                                      utils::VectorRef<const ast::Attribute*> attributes,
+                                     std::optional<uint32_t> location,
                                      std::vector<StageVariable>& variables) const;
 
     /// Recursively determine if the type contains builtin.
@@ -230,6 +236,10 @@
     /// whenever a set of expressions are resolved to globals.
     template <size_t N, typename F>
     void GetOriginatingResources(std::array<const ast::Expression*, N> exprs, F&& cb);
+
+    /// @param func the function of the entry point. Must be non-nullptr and true for IsEntryPoint()
+    /// @returns the entry point information
+    EntryPoint GetEntryPoint(const tint::ast::Function* func);
 };
 
 }  // namespace tint::inspector
diff --git a/src/tint/inspector/inspector_test.cc b/src/tint/inspector/inspector_test.cc
index 5bfd08e..5a190f6 100644
--- a/src/tint/inspector/inspector_test.cc
+++ b/src/tint/inspector/inspector_test.cc
@@ -291,7 +291,7 @@
 
     auto* in_var = Param("in_var", tint_type(),
                          utils::Vector{
-                             Location(0u),
+                             Location(0_u),
                              Flat(),
                          });
     Func("foo", utils::Vector{in_var}, tint_type(),
@@ -302,7 +302,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0u),
+             Location(0_u),
          });
     Inspector& inspector = Build();
 
@@ -336,17 +336,17 @@
 TEST_F(InspectorGetEntryPointTest, MultipleInOutVariables) {
     auto* in_var0 = Param("in_var0", ty.u32(),
                           utils::Vector{
-                              Location(0u),
+                              Location(0_u),
                               Flat(),
                           });
     auto* in_var1 = Param("in_var1", ty.u32(),
                           utils::Vector{
-                              Location(1u),
+                              Location(1_u),
                               Flat(),
                           });
     auto* in_var4 = Param("in_var4", ty.u32(),
                           utils::Vector{
-                              Location(4u),
+                              Location(4_u),
                               Flat(),
                           });
     Func("foo", utils::Vector{in_var0, in_var1, in_var4}, ty.u32(),
@@ -357,7 +357,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0u),
+             Location(0_u),
          });
     Inspector& inspector = Build();
 
@@ -393,7 +393,7 @@
 TEST_F(InspectorGetEntryPointTest, MultipleEntryPointsInOutVariables) {
     auto* in_var_foo = Param("in_var_foo", ty.u32(),
                              utils::Vector{
-                                 Location(0u),
+                                 Location(0_u),
                                  Flat(),
                              });
     Func("foo", utils::Vector{in_var_foo}, ty.u32(),
@@ -404,12 +404,12 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0u),
+             Location(0_u),
          });
 
     auto* in_var_bar = Param("in_var_bar", ty.u32(),
                              utils::Vector{
-                                 Location(0u),
+                                 Location(0_u),
                                  Flat(),
                              });
     Func("bar", utils::Vector{in_var_bar}, ty.u32(),
@@ -420,7 +420,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(1u),
+             Location(1_u),
          });
 
     Inspector& inspector = Build();
@@ -464,7 +464,7 @@
                           });
     auto* in_var1 = Param("in_var1", ty.f32(),
                           utils::Vector{
-                              Location(0u),
+                              Location(0_u),
                           });
     Func("foo", utils::Vector{in_var0, in_var1}, ty.f32(),
          utils::Vector{
@@ -596,8 +596,8 @@
          utils::Vector{
              Param("param_a", ty.Of(struct_a)),
              Param("param_b", ty.Of(struct_b)),
-             Param("param_c", ty.f32(), utils::Vector{Location(3u)}),
-             Param("param_d", ty.f32(), utils::Vector{Location(4u)}),
+             Param("param_c", ty.f32(), utils::Vector{Location(3_u)}),
+             Param("param_d", ty.f32(), utils::Vector{Location(4_u)}),
          },
          ty.Of(struct_a),
          utils::Vector{
@@ -1136,7 +1136,7 @@
 
 TEST_F(InspectorGetEntryPointTest, ImplicitInterpolate) {
     Structure("in_struct", utils::Vector{
-                               Member("struct_inner", ty.f32(), utils::Vector{Location(0)}),
+                               Member("struct_inner", ty.f32(), utils::Vector{Location(0_a)}),
                            });
 
     Func("ep_func",
@@ -1167,7 +1167,7 @@
         "in_struct",
         utils::Vector{
             Member("struct_inner", ty.f32(),
-                   utils::Vector{Interpolate(params.in_type, params.in_sampling), Location(0)}),
+                   utils::Vector{Interpolate(params.in_type, params.in_sampling), Location(0_a)}),
         });
 
     Func("ep_func",
diff --git a/src/tint/inspector/test_inspector_builder.cc b/src/tint/inspector/test_inspector_builder.cc
index 342167c..ce341a6 100644
--- a/src/tint/inspector/test_inspector_builder.cc
+++ b/src/tint/inspector/test_inspector_builder.cc
@@ -54,7 +54,7 @@
         std::tie(member_name, location) = var;
         members.Push(Member(member_name, ty.u32(),
                             utils::Vector{
-                                Location(location),
+                                Location(AInt(location)),
                                 Flat(),
                             }));
     }
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 8a80bf5..39e0b14 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -934,23 +934,23 @@
 op && (bool, bool) -> bool
 op || (bool, bool) -> bool
 
-op == <T: scalar>(T, T) -> bool
-op == <T: scalar, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op == <T: abstract_or_scalar>(T, T) -> bool
+@const op == <T: abstract_or_scalar, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
-op != <T: scalar>(T, T) -> bool
-op != <T: scalar, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op != <T: abstract_or_scalar>(T, T) -> bool
+@const op != <T: abstract_or_scalar, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
-op < <T: fiu32_f16>(T, T) -> bool
-op < <T: fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op < <T: fia_fiu32_f16>(T, T) -> bool
+@const op < <T: fia_fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
-op > <T: fiu32_f16>(T, T) -> bool
-op > <T: fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op > <T: fia_fiu32_f16>(T, T) -> bool
+@const op > <T: fia_fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
-op <= <T: fiu32_f16>(T, T) -> bool
-op <= <T: fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op <= <T: fia_fiu32_f16>(T, T) -> bool
+@const op <= <T: fia_fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
-op >= <T: fiu32_f16>(T, T) -> bool
-op >= <T: fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
+@const op >= <T: fia_fiu32_f16>(T, T) -> bool
+@const op >= <T: fiu32_f16, N: num> (vec<N, T>, vec<N, T>) -> vec<N, bool>
 
 op << <T: iu32>(T, u32) -> T
 op << <T: iu32, N: num> (vec<N, T>, vec<N, u32>) -> vec<N, T>
diff --git a/src/tint/number.h b/src/tint/number.h
index 29ae227..6e032d7 100644
--- a/src/tint/number.h
+++ b/src/tint/number.h
@@ -260,7 +260,7 @@
 using f16 = Number<detail::NumberKindF16>;
 
 /// @returns the friendly name of Number type T
-template <typename T, typename = traits::EnableIf<IsNumber<T>>>
+template <typename T, traits::EnableIf<IsNumber<T>>* = nullptr>
 const char* FriendlyName() {
     if constexpr (std::is_same_v<T, AInt>) {
         return "abstract-int";
@@ -279,6 +279,12 @@
     }
 }
 
+/// @returns the friendly name of T when T is bool
+template <typename T, traits::EnableIf<std::is_same_v<T, bool>>* = nullptr>
+const char* FriendlyName() {
+    return "bool";
+}
+
 /// Enumerator of failure reasons when converting from one number to another.
 enum class ConversionFailure {
     kExceedsPositiveLimit,  // The value was too big (+'ve) to fit in the target type
diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h
index bdd2807..ead13e9 100644
--- a/src/tint/program_builder.h
+++ b/src/tint/program_builder.h
@@ -2928,17 +2928,19 @@
 
     /// Creates an ast::LocationAttribute
     /// @param source the source information
-    /// @param location the location value
+    /// @param location the location value expression
     /// @returns the location attribute pointer
-    const ast::LocationAttribute* Location(const Source& source, uint32_t location) {
-        return create<ast::LocationAttribute>(source, location);
+    template <typename EXPR>
+    const ast::LocationAttribute* Location(const Source& source, EXPR&& location) {
+        return create<ast::LocationAttribute>(source, Expr(std::forward<EXPR>(location)));
     }
 
     /// Creates an ast::LocationAttribute
-    /// @param location the location value
+    /// @param location the location value expression
     /// @returns the location attribute pointer
-    const ast::LocationAttribute* Location(uint32_t location) {
-        return create<ast::LocationAttribute>(source_, location);
+    template <typename EXPR>
+    const ast::LocationAttribute* Location(EXPR&& location) {
+        return create<ast::LocationAttribute>(source_, Expr(std::forward<EXPR>(location)));
     }
 
     /// Creates an ast::IdAttribute
diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc
index 0acfa1d..773ba46 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -1109,7 +1109,8 @@
             // Replace this location attribute with a new one with one higher index.
             // The old one doesn't leak because it's kept in the builder's AST node
             // list.
-            attr = builder_.Location(loc_attr->source, loc_attr->value + 1);
+            attr = builder_.Location(
+                loc_attr->source, AInt(loc_attr->expr->As<ast::IntLiteralExpression>()->value + 1));
         }
     }
 }
diff --git a/src/tint/reader/spirv/parser_impl.cc b/src/tint/reader/spirv/parser_impl.cc
index 2942591..37ea5ca 100644
--- a/src/tint/reader/spirv/parser_impl.cc
+++ b/src/tint/reader/spirv/parser_impl.cc
@@ -1723,25 +1723,22 @@
     return result;
 }
 
-const ast::Attribute* ParserImpl::SetLocation(AttributeList* attributes,
-                                              const ast::Attribute* replacement) {
+void ParserImpl::SetLocation(AttributeList* attributes, const ast::Attribute* replacement) {
     if (!replacement) {
-        return nullptr;
+        return;
     }
     for (auto*& attribute : *attributes) {
         if (attribute->Is<ast::LocationAttribute>()) {
             // Replace this location attribute with the replacement.
             // The old one doesn't leak because it's kept in the builder's AST node
             // list.
-            const ast::Attribute* result = nullptr;
-            result = attribute;
             attribute = replacement;
-            return result;  // Assume there is only one such decoration.
+            return;  // Assume there is only one such decoration.
         }
     }
     // The list didn't have a location. Add it.
     attributes->Push(replacement);
-    return nullptr;
+    return;
 }
 
 bool ParserImpl::ConvertPipelineDecorations(const Type* store_type,
@@ -1759,7 +1756,7 @@
                     return Fail() << "malformed Location decoration on ID requires one "
                                      "literal operand";
                 }
-                SetLocation(attributes, create<ast::LocationAttribute>(Source{}, deco[1]));
+                SetLocation(attributes, builder_.Location(AInt(deco[1])));
                 if (store_type->IsIntegerScalarOrVector()) {
                     // Default to flat interpolation for integral user-defined IO types.
                     type = ast::InterpolationType::kFlat;
diff --git a/src/tint/reader/spirv/parser_impl.h b/src/tint/reader/spirv/parser_impl.h
index 12d6226..948a9a8 100644
--- a/src/tint/reader/spirv/parser_impl.h
+++ b/src/tint/reader/spirv/parser_impl.h
@@ -280,9 +280,7 @@
     /// Assumes the list contains at most one Location decoration.
     /// @param decos the attribute list to modify
     /// @param replacement the location decoration to place into the list
-    /// @returns the location decoration that was replaced, if one was replaced,
-    /// or null otherwise.
-    const ast::Attribute* SetLocation(AttributeList* decos, const ast::Attribute* replacement);
+    void SetLocation(AttributeList* decos, const ast::Attribute* replacement);
 
     /// Converts a SPIR-V struct member decoration into a number of AST
     /// decorations. If the decoration is recognized but deliberately dropped,
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index d57ad78..f3b424e 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -3551,7 +3551,9 @@
             }
             match(Token::Type::kComma);
 
-            return create<ast::LocationAttribute>(t.source(), val.value);
+            return builder_.Location(t.source(),
+                                     create<ast::IntLiteralExpression>(
+                                         val.value, ast::IntLiteralExpression::Suffix::kNone));
         });
     }
 
diff --git a/src/tint/reader/wgsl/parser_impl_function_decl_test.cc b/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
index cba4a7d..c9f55cf 100644
--- a/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
@@ -256,7 +256,10 @@
     ASSERT_EQ(ret_type_attributes.Length(), 1u);
     auto* loc = ret_type_attributes[0]->As<ast::LocationAttribute>();
     ASSERT_TRUE(loc != nullptr);
-    EXPECT_EQ(loc->value, 1u);
+    EXPECT_TRUE(loc->expr->Is<ast::IntLiteralExpression>());
+
+    auto* exp = loc->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(1u, exp->value);
 
     auto* body = f->body;
     ASSERT_EQ(body->statements.Length(), 1u);
diff --git a/src/tint/reader/wgsl/parser_impl_function_header_test.cc b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
index 1a8704e..c65c6e0 100644
--- a/src/tint/reader/wgsl/parser_impl_function_header_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
@@ -54,9 +54,12 @@
     EXPECT_EQ(f->params.Length(), 0u);
     EXPECT_TRUE(f->return_type->Is<ast::F32>());
     ASSERT_EQ(f->return_type_attributes.Length(), 1u);
+
     auto* loc = f->return_type_attributes[0]->As<ast::LocationAttribute>();
     ASSERT_TRUE(loc != nullptr);
-    EXPECT_EQ(loc->value, 1u);
+    ASSERT_TRUE(loc->expr->Is<ast::IntLiteralExpression>());
+    auto* exp = loc->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(exp->value, 1u);
 }
 
 TEST_F(ParserImplTest, FunctionHeader_InvariantReturnType) {
diff --git a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
index 2f1397f..01ed4d9 100644
--- a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
@@ -201,7 +201,7 @@
 
     auto* override_attr = ast::GetAttribute<ast::IdAttribute>(override->attributes);
     ASSERT_NE(override_attr, nullptr);
-    EXPECT_TRUE(override_attr->value->Is<ast::IntLiteralExpression>());
+    EXPECT_TRUE(override_attr->expr->Is<ast::IntLiteralExpression>());
 }
 
 TEST_F(ParserImplTest, GlobalOverrideDecl_WithId_TrailingComma) {
@@ -231,7 +231,7 @@
 
     auto* override_attr = ast::GetAttribute<ast::IdAttribute>(override->attributes);
     ASSERT_NE(override_attr, nullptr);
-    EXPECT_TRUE(override_attr->value->Is<ast::IntLiteralExpression>());
+    EXPECT_TRUE(override_attr->expr->Is<ast::IntLiteralExpression>());
 }
 
 TEST_F(ParserImplTest, GlobalOverrideDecl_WithoutId) {
diff --git a/src/tint/reader/wgsl/parser_impl_param_list_test.cc b/src/tint/reader/wgsl/parser_impl_param_list_test.cc
index ce542e4..26c1852 100644
--- a/src/tint/reader/wgsl/parser_impl_param_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_param_list_test.cc
@@ -117,8 +117,12 @@
     EXPECT_TRUE(e.value[1]->Is<ast::Parameter>());
     auto attrs_1 = e.value[1]->attributes;
     ASSERT_EQ(attrs_1.Length(), 1u);
-    EXPECT_TRUE(attrs_1[0]->Is<ast::LocationAttribute>());
-    EXPECT_EQ(attrs_1[0]->As<ast::LocationAttribute>()->value, 1u);
+
+    ASSERT_TRUE(attrs_1[0]->Is<ast::LocationAttribute>());
+    auto* attr = attrs_1[0]->As<ast::LocationAttribute>();
+    ASSERT_TRUE(attr->expr->Is<ast::IntLiteralExpression>());
+    auto* loc = attr->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(loc->value, 1u);
 
     EXPECT_EQ(e.value[1]->source.range.begin.line, 1u);
     EXPECT_EQ(e.value[1]->source.range.begin.column, 52u);
diff --git a/src/tint/reader/wgsl/parser_impl_struct_member_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_struct_member_attribute_test.cc
index 4fb9528..cc6d186 100644
--- a/src/tint/reader/wgsl/parser_impl_struct_member_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_struct_member_attribute_test.cc
@@ -102,9 +102,9 @@
     ASSERT_TRUE(member_attr->Is<ast::StructMemberAlignAttribute>());
 
     auto* o = member_attr->As<ast::StructMemberAlignAttribute>();
-    ASSERT_TRUE(o->align->Is<ast::IntLiteralExpression>());
-    EXPECT_EQ(o->align->As<ast::IntLiteralExpression>()->value, 4);
-    EXPECT_EQ(o->align->As<ast::IntLiteralExpression>()->suffix,
+    ASSERT_TRUE(o->expr->Is<ast::IntLiteralExpression>());
+    EXPECT_EQ(o->expr->As<ast::IntLiteralExpression>()->value, 4);
+    EXPECT_EQ(o->expr->As<ast::IntLiteralExpression>()->suffix,
               ast::IntLiteralExpression::Suffix::kNone);
 }
 
@@ -121,9 +121,9 @@
     ASSERT_TRUE(member_attr->Is<ast::StructMemberAlignAttribute>());
 
     auto* o = member_attr->As<ast::StructMemberAlignAttribute>();
-    ASSERT_TRUE(o->align->Is<ast::IntLiteralExpression>());
+    ASSERT_TRUE(o->expr->Is<ast::IntLiteralExpression>());
 
-    auto* expr = o->align->As<ast::IntLiteralExpression>();
+    auto* expr = o->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 }
diff --git a/src/tint/reader/wgsl/parser_impl_struct_member_test.cc b/src/tint/reader/wgsl/parser_impl_struct_member_test.cc
index c64f772..fe233fe 100644
--- a/src/tint/reader/wgsl/parser_impl_struct_member_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_struct_member_test.cc
@@ -51,9 +51,8 @@
     EXPECT_TRUE(m->attributes[0]->Is<ast::StructMemberAlignAttribute>());
 
     auto* attr = m->attributes[0]->As<ast::StructMemberAlignAttribute>();
-    ASSERT_TRUE(attr->align->Is<ast::IntLiteralExpression>());
-
-    auto* expr = attr->align->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(attr->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = attr->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 2);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 
@@ -101,8 +100,8 @@
     ASSERT_TRUE(m->attributes[1]->Is<ast::StructMemberAlignAttribute>());
     auto* attr = m->attributes[1]->As<ast::StructMemberAlignAttribute>();
 
-    ASSERT_TRUE(attr->align->Is<ast::IntLiteralExpression>());
-    auto* expr = attr->align->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(attr->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = attr->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 
diff --git a/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc b/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
index 2745e5f..a814ea6 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
@@ -31,7 +31,12 @@
     ASSERT_NE(attr_1, nullptr);
 
     ASSERT_TRUE(attr_0->Is<ast::LocationAttribute>());
-    EXPECT_EQ(attr_0->As<ast::LocationAttribute>()->value, 4u);
+
+    auto* loc = attr_0->As<ast::LocationAttribute>();
+    ASSERT_TRUE(loc->expr->Is<ast::IntLiteralExpression>());
+    auto* exp = loc->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(exp->value, 4u);
+
     ASSERT_TRUE(attr_1->Is<ast::BuiltinAttribute>());
     EXPECT_EQ(attr_1->As<ast::BuiltinAttribute>()->builtin, ast::BuiltinValue::kPosition);
 }
diff --git a/src/tint/reader/wgsl/parser_impl_variable_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_variable_attribute_test.cc
index e48cc85..7c52838 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_attribute_test.cc
@@ -29,7 +29,9 @@
     ASSERT_TRUE(var_attr->Is<ast::LocationAttribute>());
 
     auto* loc = var_attr->As<ast::LocationAttribute>();
-    EXPECT_EQ(loc->value, 4u);
+    ASSERT_TRUE(loc->expr->Is<ast::IntLiteralExpression>());
+    auto* exp = loc->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(exp->value, 4u);
 }
 
 TEST_F(ParserImplTest, Attribute_Location_TrailingComma) {
@@ -44,7 +46,9 @@
     ASSERT_TRUE(var_attr->Is<ast::LocationAttribute>());
 
     auto* loc = var_attr->As<ast::LocationAttribute>();
-    EXPECT_EQ(loc->value, 4u);
+    ASSERT_TRUE(loc->expr->Is<ast::IntLiteralExpression>());
+    auto* exp = loc->expr->As<ast::IntLiteralExpression>();
+    EXPECT_EQ(exp->value, 4u);
 }
 
 TEST_F(ParserImplTest, Attribute_Location_MissingLeftParen) {
@@ -364,9 +368,8 @@
     ASSERT_TRUE(var_attr->Is<ast::BindingAttribute>());
 
     auto* binding = var_attr->As<ast::BindingAttribute>();
-    ASSERT_TRUE(binding->value->Is<ast::IntLiteralExpression>());
-
-    auto* expr = binding->value->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(binding->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = binding->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 }
@@ -383,9 +386,8 @@
     ASSERT_TRUE(var_attr->Is<ast::BindingAttribute>());
 
     auto* binding = var_attr->As<ast::BindingAttribute>();
-    ASSERT_TRUE(binding->value->Is<ast::IntLiteralExpression>());
-
-    auto* expr = binding->value->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(binding->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = binding->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 }
@@ -442,9 +444,8 @@
     ASSERT_TRUE(var_attr->Is<ast::GroupAttribute>());
 
     auto* group = var_attr->As<ast::GroupAttribute>();
-    ASSERT_TRUE(group->value->Is<ast::IntLiteralExpression>());
-
-    auto* expr = group->value->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(group->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = group->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 }
@@ -461,9 +462,8 @@
     ASSERT_TRUE(var_attr->Is<ast::GroupAttribute>());
 
     auto* group = var_attr->As<ast::GroupAttribute>();
-    ASSERT_TRUE(group->value->Is<ast::IntLiteralExpression>());
-
-    auto* expr = group->value->As<ast::IntLiteralExpression>();
+    ASSERT_TRUE(group->expr->Is<ast::IntLiteralExpression>());
+    auto* expr = group->expr->As<ast::IntLiteralExpression>();
     EXPECT_EQ(expr->value, 4);
     EXPECT_EQ(expr->suffix, ast::IntLiteralExpression::Suffix::kNone);
 }
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index 7f573a8..b182f6d 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -104,7 +104,7 @@
         case AttributeKind::kInvariant:
             return {builder.Invariant(source)};
         case AttributeKind::kLocation:
-            return {builder.Location(source, 1)};
+            return {builder.Location(source, 1_a)};
         case AttributeKind::kOffset:
             return {builder.create<ast::StructMemberOffsetAttribute>(source, 4u)};
         case AttributeKind::kSize:
@@ -286,7 +286,7 @@
     auto& params = GetParam();
     auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
     if (params.kind != AttributeKind::kLocation) {
-        attrs.Push(Location(Source{{34, 56}}, 2));
+        attrs.Push(Location(Source{{34, 56}}, 2_a));
     }
     auto* p = Param("a", ty.vec4<f32>(), attrs);
     Func("vertex_main", utils::Vector{p}, ty.vec4<f32>(),
@@ -388,7 +388,7 @@
 TEST_P(FragmentShaderReturnTypeAttributeTest, IsValid) {
     auto& params = GetParam();
     auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
-    attrs.Push(Location(Source{{34, 56}}, 2));
+    attrs.Push(Location(Source{{34, 56}}, 2_a));
     Func("frag_main", utils::Empty, ty.vec4<f32>(),
          utils::Vector{Return(Construct(ty.vec4<f32>()))},
          utils::Vector{
@@ -495,8 +495,8 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(Source{{12, 34}}, 2),
-             Location(Source{{56, 78}}, 3),
+             Location(Source{{12, 34}}, 2_a),
+             Location(Source{{56, 78}}, 3_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -531,8 +531,8 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(Source{{12, 34}}, 2),
-             Location(Source{{56, 78}}, 3),
+             Location(Source{{12, 34}}, 2_a),
+             Location(Source{{56, 78}}, 3_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -698,12 +698,12 @@
 TEST_P(VariableAttributeTest, IsValid) {
     auto& params = GetParam();
 
+    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
+    auto* attr = attrs[0];
     if (IsBindingAttribute(params.kind)) {
-        GlobalVar("a", ty.sampler(ast::SamplerKind::kSampler),
-                  createAttributes(Source{{12, 34}}, *this, params.kind));
+        GlobalVar("a", ty.sampler(ast::SamplerKind::kSampler), attrs);
     } else {
-        GlobalVar("a", ty.f32(), ast::StorageClass::kPrivate,
-                  createAttributes(Source{{12, 34}}, *this, params.kind));
+        GlobalVar("a", ty.f32(), ast::StorageClass::kPrivate, attrs);
     }
 
     if (params.should_pass) {
@@ -711,7 +711,8 @@
     } else {
         EXPECT_FALSE(r()->Resolve());
         if (!IsBindingAttribute(params.kind)) {
-            EXPECT_EQ(r()->error(), "12:34 error: attribute is not valid for module-scope 'var'");
+            EXPECT_EQ(r()->error(), "12:34 error: attribute '" + attr->Name() +
+                                        "' is not valid for module-scope 'var'");
         }
     }
 }
@@ -1100,7 +1101,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
@@ -1109,7 +1110,7 @@
     auto* param = Param("p", ty.vec4<f32>(),
                         utils::Vector{
                             Invariant(Source{{12, 34}}),
-                            Location(0),
+                            Location(0_a),
                         });
     Func("main", utils::Vector{param}, ty.vec4<f32>(),
          utils::Vector{
@@ -1119,7 +1120,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
@@ -1218,7 +1219,7 @@
          utils::Vector{
              Param("a", ty.f32(),
                    utils::Vector{
-                       Location(0),
+                       Location(0_a),
                        Interpolate(Source{{12, 34}}, params.type, params.sampling),
                    }),
          },
@@ -1244,7 +1245,7 @@
          utils::Vector{
              Param("a", ty.i32(),
                    utils::Vector{
-                       Location(0),
+                       Location(0_a),
                        Interpolate(Source{{12, 34}}, params.type, params.sampling),
                    }),
          },
@@ -1275,7 +1276,7 @@
          utils::Vector{
              Param("a", ty.vec4<u32>(),
                    utils::Vector{
-                       Location(0),
+                       Location(0_a),
                        Interpolate(Source{{12, 34}}, params.type, params.sampling),
                    }),
          },
@@ -1318,7 +1319,8 @@
         Params{ast::InterpolationType::kFlat, ast::InterpolationSampling::kSample, false}));
 
 TEST_F(InterpolateTest, FragmentInput_Integer_MissingFlatInterpolation) {
-    Func("main", utils::Vector{Param(Source{{12, 34}}, "a", ty.i32(), utils::Vector{Location(0)})},
+    Func("main",
+         utils::Vector{Param(Source{{12, 34}}, "a", ty.i32(), utils::Vector{Location(0_a)})},
          ty.void_(), utils::Empty,
          utils::Vector{
              Stage(ast::PipelineStage::kFragment),
@@ -1335,7 +1337,7 @@
         "S",
         utils::Vector{
             Member("pos", ty.vec4<f32>(), utils::Vector{Builtin(ast::BuiltinValue::kPosition)}),
-            Member(Source{{12, 34}}, "u", ty.u32(), utils::Vector{Location(0)}),
+            Member(Source{{12, 34}}, "u", ty.u32(), utils::Vector{Location(0_a)}),
         });
     Func("main", utils::Empty, ty.Of(s),
          utils::Vector{
diff --git a/src/tint/resolver/builtins_validation_test.cc b/src/tint/resolver/builtins_validation_test.cc
index 6318901..a27052c 100644
--- a/src/tint/resolver/builtins_validation_test.cc
+++ b/src/tint/resolver/builtins_validation_test.cc
@@ -163,7 +163,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
@@ -198,7 +198,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
@@ -256,7 +256,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -301,7 +301,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -330,7 +330,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -372,7 +372,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(), "12:34 error: store type of builtin(sample_mask) must be 'u32'");
@@ -400,7 +400,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -427,7 +427,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(), "12:34 error: store type of builtin(sample_index) must be 'u32'");
@@ -453,7 +453,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(), "12:34 error: store type of builtin(position) must be 'vec4<f32>'");
@@ -745,7 +745,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
@@ -768,7 +768,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -797,7 +797,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 0f4872b..991fe79 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -94,6 +94,21 @@
 /// Helper that calls `f` passing in the value of all `cs`.
 /// Assumes all `cs` are of the same type.
 template <typename F, typename... CONSTANTS>
+auto Dispatch_fia_fiu32_f16_bool(F&& f, CONSTANTS&&... cs) {
+    return Switch(
+        First(cs...)->Type(),  //
+        [&](const sem::AbstractInt*) { return f(cs->template As<AInt>()...); },
+        [&](const sem::AbstractFloat*) { return f(cs->template As<AFloat>()...); },
+        [&](const sem::F32*) { return f(cs->template As<f32>()...); },
+        [&](const sem::I32*) { return f(cs->template As<i32>()...); },
+        [&](const sem::U32*) { return f(cs->template As<u32>()...); },
+        [&](const sem::F16*) { return f(cs->template As<f16>()...); },
+        [&](const sem::Bool*) { return f(cs->template As<bool>()...); });
+}
+
+/// Helper that calls `f` passing in the value of all `cs`.
+/// Assumes all `cs` are of the same type.
+template <typename F, typename... CONSTANTS>
 auto Dispatch_fa_f32_f16(F&& f, CONSTANTS&&... cs) {
     return Switch(
         First(cs...)->Type(),  //
@@ -466,10 +481,14 @@
     }
 }
 
-/// TransformElements constructs a new constant by applying the transformation function 'f' on each
-/// of the most deeply nested elements of 'cs'. Assumes that all constants are the same type.
+/// TransformElements constructs a new constant of type `composite_ty` by applying the
+/// transformation function 'f' on each of the most deeply nested elements of 'cs'. Assumes that all
+/// input constants `cs` are of the same type.
 template <typename F, typename... CONSTANTS>
-const Constant* TransformElements(ProgramBuilder& builder, F&& f, CONSTANTS&&... cs) {
+const Constant* TransformElements(ProgramBuilder& builder,
+                                  const sem::Type* composite_ty,
+                                  F&& f,
+                                  CONSTANTS&&... cs) {
     uint32_t n = 0;
     auto* ty = First(cs...)->Type();
     auto* el_ty = sem::Type::ElementOf(ty, &n);
@@ -479,16 +498,19 @@
     utils::Vector<const sem::Constant*, 8> els;
     els.Reserve(n);
     for (uint32_t i = 0; i < n; i++) {
-        els.Push(TransformElements(builder, std::forward<F>(f), cs->Index(i)...));
+        els.Push(TransformElements(builder, sem::Type::ElementOf(composite_ty), std::forward<F>(f),
+                                   cs->Index(i)...));
     }
-    return CreateComposite(builder, ty, std::move(els));
+    return CreateComposite(builder, composite_ty, std::move(els));
 }
 
-/// TransformBinaryElements constructs a new constant by applying the transformation function 'f' on
-/// each of the most deeply nested elements of both `c0` and `c1`. Unlike TransformElements, this
-/// function handles the constants being of different types, e.g. vector-scalar, scalar-vector.
+/// TransformBinaryElements constructs a new constant of type `composite_ty` by applying the
+/// transformation function 'f' on each of the most deeply nested elements of both `c0` and `c1`.
+/// Unlike TransformElements, this function handles the constants being of different types, e.g.
+/// vector-scalar, scalar-vector.
 template <typename F>
 const Constant* TransformBinaryElements(ProgramBuilder& builder,
+                                        const sem::Type* composite_ty,
                                         F&& f,
                                         const sem::Constant* c0,
                                         const sem::Constant* c1) {
@@ -510,12 +532,11 @@
             }
             return c->Index(i);
         };
-        els.Push(TransformBinaryElements(builder, std::forward<F>(f), nested_or_self(c0, n0),
+        els.Push(TransformBinaryElements(builder, sem::Type::ElementOf(composite_ty),
+                                         std::forward<F>(f), nested_or_self(c0, n0),
                                          nested_or_self(c1, n1)));
     }
-    // Use larger type
-    auto* ty = n0 > n1 ? c0->Type() : c1->Type();
-    return CreateComposite(builder, ty, std::move(els));
+    return CreateComposite(builder, composite_ty, std::move(els));
 }
 }  // namespace
 
@@ -915,7 +936,7 @@
     return nullptr;
 }
 
-ConstEval::ConstantResult ConstEval::OpComplement(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpComplement(const sem::Type* ty,
                                                   utils::VectorRef<const sem::Constant*> args,
                                                   const Source&) {
     auto transform = [&](const sem::Constant* c) {
@@ -924,10 +945,10 @@
         };
         return Dispatch_ia_iu32(create, c);
     };
-    return TransformElements(builder, transform, args[0]);
+    return TransformElements(builder, ty, transform, args[0]);
 }
 
-ConstEval::ConstantResult ConstEval::OpUnaryMinus(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpUnaryMinus(const sem::Type* ty,
                                                   utils::VectorRef<const sem::Constant*> args,
                                                   const Source&) {
     auto transform = [&](const sem::Constant* c) {
@@ -949,10 +970,10 @@
         };
         return Dispatch_fia_fi32_f16(create, c);
     };
-    return TransformElements(builder, transform, args[0]);
+    return TransformElements(builder, ty, transform, args[0]);
 }
 
-ConstEval::ConstantResult ConstEval::OpPlus(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpPlus(const sem::Type* ty,
                                             utils::VectorRef<const sem::Constant*> args,
                                             const Source& source) {
     TINT_SCOPED_ASSIGNMENT(current_source, &source);
@@ -963,14 +984,14 @@
         return nullptr;
     };
 
-    auto r = TransformBinaryElements(builder, transform, args[0], args[1]);
+    auto r = TransformBinaryElements(builder, ty, transform, args[0], args[1]);
     if (builder.Diagnostics().contains_errors()) {
         return utils::Failure;
     }
     return r;
 }
 
-ConstEval::ConstantResult ConstEval::OpMinus(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpMinus(const sem::Type* ty,
                                              utils::VectorRef<const sem::Constant*> args,
                                              const Source& source) {
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
@@ -1003,14 +1024,14 @@
         return Dispatch_fia_fiu32_f16(create, c0, c1);
     };
 
-    auto r = TransformBinaryElements(builder, transform, args[0], args[1]);
+    auto r = TransformBinaryElements(builder, ty, transform, args[0], args[1]);
     if (builder.Diagnostics().contains_errors()) {
         return utils::Failure;
     }
     return r;
 }
 
-ConstEval::ConstantResult ConstEval::OpMultiply(const sem::Type* /*ty*/,
+ConstEval::ConstantResult ConstEval::OpMultiply(const sem::Type* ty,
                                                 utils::VectorRef<const sem::Constant*> args,
                                                 const Source& source) {
     TINT_SCOPED_ASSIGNMENT(current_source, &source);
@@ -1021,7 +1042,7 @@
         return nullptr;
     };
 
-    auto r = TransformBinaryElements(builder, transform, args[0], args[1]);
+    auto r = TransformBinaryElements(builder, ty, transform, args[0], args[1]);
     if (builder.Diagnostics().contains_errors()) {
         return utils::Failure;
     }
@@ -1196,7 +1217,7 @@
     return CreateComposite(builder, ty, result_mat);
 }
 
-ConstEval::ConstantResult ConstEval::OpDivide(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpDivide(const sem::Type* ty,
                                               utils::VectorRef<const sem::Constant*> args,
                                               const Source& source) {
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
@@ -1237,14 +1258,116 @@
         return Dispatch_fia_fiu32_f16(create, c0, c1);
     };
 
-    auto r = TransformBinaryElements(builder, transform, args[0], args[1]);
+    auto r = TransformBinaryElements(builder, ty, transform, args[0], args[1]);
     if (builder.Diagnostics().contains_errors()) {
         return utils::Failure;
     }
     return r;
 }
 
-ConstEval::ConstantResult ConstEval::atan2(const sem::Type*,
+ConstEval::ConstantResult ConstEval::OpEqual(const sem::Type* ty,
+                                             utils::VectorRef<const sem::Constant*> args,
+                                             const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i == j);
+        };
+        return Dispatch_fia_fiu32_f16_bool(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::OpNotEqual(const sem::Type* ty,
+                                                utils::VectorRef<const sem::Constant*> args,
+                                                const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i != j);
+        };
+        return Dispatch_fia_fiu32_f16_bool(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::OpLessThan(const sem::Type* ty,
+                                                utils::VectorRef<const sem::Constant*> args,
+                                                const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i < j);
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::OpGreaterThan(const sem::Type* ty,
+                                                   utils::VectorRef<const sem::Constant*> args,
+                                                   const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i > j);
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::OpLessThanEqual(const sem::Type* ty,
+                                                     utils::VectorRef<const sem::Constant*> args,
+                                                     const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i <= j);
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::OpGreaterThanEqual(const sem::Type* ty,
+                                                        utils::VectorRef<const sem::Constant*> args,
+                                                        const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto i, auto j) -> const Constant* {
+            return CreateElement(builder, sem::Type::DeepestElementOf(ty), i >= j);
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+
+    auto r = TransformElements(builder, ty, transform, args[0], args[1]);
+    if (builder.Diagnostics().contains_errors()) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+ConstEval::ConstantResult ConstEval::atan2(const sem::Type* ty,
                                            utils::VectorRef<const sem::Constant*> args,
                                            const Source&) {
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
@@ -1253,10 +1376,10 @@
         };
         return Dispatch_fa_f32_f16(create, c0, c1);
     };
-    return TransformElements(builder, transform, args[0], args[1]);
+    return TransformElements(builder, ty, transform, args[0], args[1]);
 }
 
-ConstEval::ConstantResult ConstEval::clamp(const sem::Type*,
+ConstEval::ConstantResult ConstEval::clamp(const sem::Type* ty,
                                            utils::VectorRef<const sem::Constant*> args,
                                            const Source&) {
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1,
@@ -1267,7 +1390,7 @@
         };
         return Dispatch_fia_fiu32_f16(create, c0, c1, c2);
     };
-    return TransformElements(builder, transform, args[0], args[1], args[2]);
+    return TransformElements(builder, ty, transform, args[0], args[1], args[2]);
 }
 
 utils::Result<const sem::Constant*> ConstEval::Convert(const sem::Type* target_ty,
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index f84e28c..df98d58 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -275,6 +275,60 @@
                             utils::VectorRef<const sem::Constant*> args,
                             const Source& source);
 
+    /// Equality operator '=='
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpEqual(const sem::Type* ty,
+                           utils::VectorRef<const sem::Constant*> args,
+                           const Source& source);
+
+    /// Inequality operator '!='
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpNotEqual(const sem::Type* ty,
+                              utils::VectorRef<const sem::Constant*> args,
+                              const Source& source);
+
+    /// Less than operator '<'
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpLessThan(const sem::Type* ty,
+                              utils::VectorRef<const sem::Constant*> args,
+                              const Source& source);
+
+    /// Greater than operator '>'
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpGreaterThan(const sem::Type* ty,
+                                 utils::VectorRef<const sem::Constant*> args,
+                                 const Source& source);
+
+    /// Less than or equal operator '<='
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpLessThanEqual(const sem::Type* ty,
+                                   utils::VectorRef<const sem::Constant*> args,
+                                   const Source& source);
+
+    /// Greater than or equal operator '>='
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    ConstantResult OpGreaterThanEqual(const sem::Type* ty,
+                                      utils::VectorRef<const sem::Constant*> args,
+                                      const Source& source);
+
     ////////////////////////////////////////////////////////////////////////////
     // Builtins
     ////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/resolver/const_eval_test.cc b/src/tint/resolver/const_eval_test.cc
index 1f745a3..8b48e8f 100644
--- a/src/tint/resolver/const_eval_test.cc
+++ b/src/tint/resolver/const_eval_test.cc
@@ -3202,6 +3202,7 @@
                            Value<i32>,
                            Value<f32>,
                            Value<f16>,
+                           Value<bool>,
 
                            Value<builder::vec2<AInt>>,
                            Value<builder::vec2<AFloat>>,
@@ -3209,6 +3210,7 @@
                            Value<builder::vec2<i32>>,
                            Value<builder::vec2<f32>>,
                            Value<builder::vec2<f16>>,
+                           Value<builder::vec2<bool>>,
 
                            Value<builder::vec3<AInt>>,
                            Value<builder::vec3<AFloat>>,
@@ -3584,6 +3586,115 @@
                                  OpDivFloatCases<f32>(),
                                  OpDivFloatCases<f16>()))));
 
+template <typename T, bool equals>
+std::vector<Case> OpEqualCases() {
+    return {
+        C(Val(T{0}), Val(T{0}), Val(true == equals)),
+        C(Val(T{0}), Val(T{1}), Val(false == equals)),
+        C(Val(T{1}), Val(T{0}), Val(false == equals)),
+        C(Val(T{1}), Val(T{1}), Val(true == equals)),
+        C(Vec(T{0}, T{0}), Vec(T{0}, T{0}), Vec(true == equals, true == equals)),
+        C(Vec(T{1}, T{0}), Vec(T{0}, T{1}), Vec(false == equals, false == equals)),
+        C(Vec(T{1}, T{1}), Vec(T{0}, T{1}), Vec(false == equals, true == equals)),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(Equal,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kEqual),
+                             testing::ValuesIn(Concat(  //
+                                 OpEqualCases<AInt, true>(),
+                                 OpEqualCases<i32, true>(),
+                                 OpEqualCases<u32, true>(),
+                                 OpEqualCases<AFloat, true>(),
+                                 OpEqualCases<f32, true>(),
+                                 OpEqualCases<f16, true>(),
+                                 OpEqualCases<bool, true>()))));
+INSTANTIATE_TEST_SUITE_P(NotEqual,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kNotEqual),
+                             testing::ValuesIn(Concat(  //
+                                 OpEqualCases<AInt, false>(),
+                                 OpEqualCases<i32, false>(),
+                                 OpEqualCases<u32, false>(),
+                                 OpEqualCases<AFloat, false>(),
+                                 OpEqualCases<f32, false>(),
+                                 OpEqualCases<f16, false>(),
+                                 OpEqualCases<bool, false>()))));
+
+template <typename T, bool less_than>
+std::vector<Case> OpLessThanCases() {
+    return {
+        C(Val(T{0}), Val(T{0}), Val(false == less_than)),
+        C(Val(T{0}), Val(T{1}), Val(true == less_than)),
+        C(Val(T{1}), Val(T{0}), Val(false == less_than)),
+        C(Val(T{1}), Val(T{1}), Val(false == less_than)),
+        C(Vec(T{0}, T{0}), Vec(T{0}, T{0}), Vec(false == less_than, false == less_than)),
+        C(Vec(T{0}, T{0}), Vec(T{1}, T{1}), Vec(true == less_than, true == less_than)),
+        C(Vec(T{1}, T{1}), Vec(T{0}, T{0}), Vec(false == less_than, false == less_than)),
+        C(Vec(T{1}, T{0}), Vec(T{0}, T{1}), Vec(false == less_than, true == less_than)),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(LessThan,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kLessThan),
+                             testing::ValuesIn(Concat(  //
+                                 OpLessThanCases<AInt, true>(),
+                                 OpLessThanCases<i32, true>(),
+                                 OpLessThanCases<u32, true>(),
+                                 OpLessThanCases<AFloat, true>(),
+                                 OpLessThanCases<f32, true>(),
+                                 OpLessThanCases<f16, true>()))));
+INSTANTIATE_TEST_SUITE_P(GreaterThanEqual,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kGreaterThanEqual),
+                             testing::ValuesIn(Concat(  //
+                                 OpLessThanCases<AInt, false>(),
+                                 OpLessThanCases<i32, false>(),
+                                 OpLessThanCases<u32, false>(),
+                                 OpLessThanCases<AFloat, false>(),
+                                 OpLessThanCases<f32, false>(),
+                                 OpLessThanCases<f16, false>()))));
+
+template <typename T, bool greater_than>
+std::vector<Case> OpGreaterThanCases() {
+    return {
+        C(Val(T{0}), Val(T{0}), Val(false == greater_than)),
+        C(Val(T{0}), Val(T{1}), Val(false == greater_than)),
+        C(Val(T{1}), Val(T{0}), Val(true == greater_than)),
+        C(Val(T{1}), Val(T{1}), Val(false == greater_than)),
+        C(Vec(T{0}, T{0}), Vec(T{0}, T{0}), Vec(false == greater_than, false == greater_than)),
+        C(Vec(T{1}, T{1}), Vec(T{0}, T{0}), Vec(true == greater_than, true == greater_than)),
+        C(Vec(T{0}, T{0}), Vec(T{1}, T{1}), Vec(false == greater_than, false == greater_than)),
+        C(Vec(T{1}, T{0}), Vec(T{0}, T{1}), Vec(true == greater_than, false == greater_than)),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(GreaterThan,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kGreaterThan),
+                             testing::ValuesIn(Concat(  //
+                                 OpGreaterThanCases<AInt, true>(),
+                                 OpGreaterThanCases<i32, true>(),
+                                 OpGreaterThanCases<u32, true>(),
+                                 OpGreaterThanCases<AFloat, true>(),
+                                 OpGreaterThanCases<f32, true>(),
+                                 OpGreaterThanCases<f16, true>()))));
+INSTANTIATE_TEST_SUITE_P(LessThanEqual,
+                         ResolverConstEvalBinaryOpTest,
+                         testing::Combine(  //
+                             testing::Values(ast::BinaryOp::kLessThanEqual),
+                             testing::ValuesIn(Concat(  //
+                                 OpGreaterThanCases<AInt, false>(),
+                                 OpGreaterThanCases<i32, false>(),
+                                 OpGreaterThanCases<u32, false>(),
+                                 OpGreaterThanCases<AFloat, false>(),
+                                 OpGreaterThanCases<f32, false>(),
+                                 OpGreaterThanCases<f16, false>()))));
+
 // Tests for errors on overflow/underflow of binary operations with abstract numbers
 struct OverflowCase {
     ast::BinaryOp op;
@@ -3608,7 +3719,7 @@
     std::string type_name = std::visit(
         [&](auto&& value) {
             using ValueType = std::decay_t<decltype(value)>;
-            return tint::FriendlyName<typename ValueType::ElementType>();
+            return builder::FriendlyName<ValueType>();
         },
         c.lhs);
 
diff --git a/src/tint/resolver/entry_point_validation_test.cc b/src/tint/resolver/entry_point_validation_test.cc
index d997912..2179f68 100644
--- a/src/tint/resolver/entry_point_validation_test.cc
+++ b/src/tint/resolver/entry_point_validation_test.cc
@@ -57,7 +57,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();
@@ -110,7 +110,7 @@
              Stage(ast::PipelineStage::kVertex),
          },
          utils::Vector{
-             Location(Source{{13, 43}}, 0),
+             Location(Source{{13, 43}}, 0_a),
              Builtin(Source{{14, 52}}, ast::BuiltinValue::kPosition),
          });
 
@@ -130,7 +130,7 @@
     // }
     auto* output = Structure(
         "Output", utils::Vector{
-                      Member("a", ty.f32(), utils::Vector{Location(0)}),
+                      Member("a", ty.f32(), utils::Vector{Location(0_a)}),
                       Member("b", ty.f32(), utils::Vector{Builtin(ast::BuiltinValue::kFragDepth)}),
                   });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
@@ -156,7 +156,7 @@
         "Output",
         utils::Vector{
             Member("a", ty.f32(),
-                   utils::Vector{Location(Source{{13, 43}}, 0),
+                   utils::Vector{Location(Source{{13, 43}}, 0_a),
                                  Builtin(Source{{14, 52}}, ast::BuiltinValue::kFragDepth)}),
         });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
@@ -182,11 +182,11 @@
     // fn main() -> Output {
     //   return Output();
     // }
-    auto* output =
-        Structure("Output", utils::Vector{
-                                Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0)}),
-                                Member(Source{{14, 52}}, "b", ty.f32(), {}),
-                            });
+    auto* output = Structure(
+        "Output", utils::Vector{
+                      Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0_a)}),
+                      Member(Source{{14, 52}}, "b", ty.f32(), {}),
+                  });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
          utils::Vector{
              Return(Construct(ty.Of(output))),
@@ -235,7 +235,7 @@
     // fn main(@location(0) param : f32) {}
     auto* param = Param("param", ty.f32(),
                         utils::Vector{
-                            Location(0),
+                            Location(0_a),
                         });
     Func(Source{{12, 34}}, "main",
          utils::Vector{
@@ -271,7 +271,7 @@
     // fn main(@location(0) @builtin(sample_index) param : u32) {}
     auto* param = Param("param", ty.u32(),
                         utils::Vector{
-                            Location(Source{{13, 43}}, 0),
+                            Location(Source{{13, 43}}, 0_a),
                             Builtin(Source{{14, 52}}, ast::BuiltinValue::kSampleIndex),
                         });
     Func(Source{{12, 34}}, "main",
@@ -297,7 +297,7 @@
     // fn main(param : Input) {}
     auto* input = Structure(
         "Input", utils::Vector{
-                     Member("a", ty.f32(), utils::Vector{Location(0)}),
+                     Member("a", ty.f32(), utils::Vector{Location(0_a)}),
                      Member("b", ty.u32(), utils::Vector{Builtin(ast::BuiltinValue::kSampleIndex)}),
                  });
     auto* param = Param("param", ty.Of(input));
@@ -323,7 +323,7 @@
         "Input",
         utils::Vector{
             Member("a", ty.u32(),
-                   utils::Vector{Location(Source{{13, 43}}, 0),
+                   utils::Vector{Location(Source{{13, 43}}, 0_a),
                                  Builtin(Source{{14, 52}}, ast::BuiltinValue::kSampleIndex)}),
         });
     auto* param = Param("param", ty.Of(input));
@@ -349,11 +349,11 @@
     // };
     // @fragment
     // fn main(param : Input) {}
-    auto* input =
-        Structure("Input", utils::Vector{
-                               Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0)}),
-                               Member(Source{{14, 52}}, "b", ty.f32(), {}),
-                           });
+    auto* input = Structure(
+        "Input", utils::Vector{
+                     Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0_a)}),
+                     Member(Source{{14, 52}}, "b", ty.f32(), {}),
+                 });
     auto* param = Param("param", ty.Of(input));
     Func(Source{{12, 34}}, "main",
          utils::Vector{
@@ -628,7 +628,7 @@
 
     auto* a = Param("a", params.create_ast_type(*this),
                     utils::Vector{
-                        Location(0),
+                        Location(0_a),
                         Flat(),
                     });
     Func(Source{{12, 34}}, "main",
@@ -657,10 +657,10 @@
 
     Enable(ast::Extension::kF16);
 
-    auto* input = Structure(
-        "Input", utils::Vector{
-                     Member("a", params.create_ast_type(*this), utils::Vector{Location(0), Flat()}),
-                 });
+    auto* input = Structure("Input", utils::Vector{
+                                         Member("a", params.create_ast_type(*this),
+                                                utils::Vector{Location(0_a), Flat()}),
+                                     });
     auto* a = Param("a", ty.Of(input), {});
     Func(Source{{12, 34}}, "main",
          utils::Vector{
@@ -695,7 +695,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     if (params.is_valid) {
@@ -719,7 +719,7 @@
 
     auto* output = Structure(
         "Output", utils::Vector{
-                      Member("a", params.create_ast_type(*this), utils::Vector{Location(0)}),
+                      Member("a", params.create_ast_type(*this), utils::Vector{Location(0_a)}),
                   });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
          utils::Vector{
@@ -751,7 +751,7 @@
 
     auto* p = Param(Source{{12, 34}}, "a", ty.i32(),
                     utils::Vector{
-                        Location(0),
+                        Location(0_a),
                         Flat(),
                     });
     Func("frag_main",
@@ -772,7 +772,7 @@
 
     auto* p = Param(Source{{12, 34}}, "a", ty.bool_(),
                     utils::Vector{
-                        Location(Source{{34, 56}}, 0),
+                        Location(Source{{34, 56}}, 0_a),
                     });
     Func("frag_main",
          utils::Vector{
@@ -803,7 +803,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(Source{{34, 56}}, 0),
+             Location(Source{{34, 56}}, 0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -825,7 +825,7 @@
                                      });
     auto* param = Param(Source{{12, 34}}, "param", ty.Of(input),
                         utils::Vector{
-                            Location(Source{{13, 43}}, 0),
+                            Location(Source{{13, 43}}, 0_a),
                         });
     Func(Source{{12, 34}}, "main",
          utils::Vector{
@@ -853,10 +853,10 @@
     // };
     // @fragment
     // fn main(param : Input) {}
-    auto* inner =
-        Structure("Inner", utils::Vector{
-                               Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0)}),
-                           });
+    auto* inner = Structure(
+        "Inner", utils::Vector{
+                     Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0_a)}),
+                 });
     auto* input = Structure("Input", utils::Vector{
                                          Member(Source{{14, 52}}, "a", ty.Of(inner)),
                                      });
@@ -884,7 +884,7 @@
     // fn main(param : Input) {}
     auto* input = Structure(
         "Input", utils::Vector{
-                     Member(Source{{13, 43}}, "a", ty.array<f32>(), utils::Vector{Location(0)}),
+                     Member(Source{{13, 43}}, "a", ty.array<f32>(), utils::Vector{Location(0_a)}),
                  });
     auto* param = Param("param", ty.Of(input));
     Func(Source{{12, 34}}, "main",
@@ -911,7 +911,7 @@
 
     auto* m = Member(Source{{34, 56}}, "m", ty.array<i32>(),
                      utils::Vector{
-                         Location(Source{{12, 34}}, 0u),
+                         Location(Source{{12, 34}}, 0_u),
                      });
     auto* s = Structure("S", utils::Vector{m});
     auto* p = Param("a", ty.Of(s));
@@ -939,7 +939,7 @@
     // fn frag_main() -> S {}
     auto* m = Member(Source{{34, 56}}, "m", ty.atomic<i32>(),
                      utils::Vector{
-                         Location(Source{{12, 34}}, 0u),
+                         Location(Source{{12, 34}}, 0_u),
                      });
     auto* s = Structure("S", utils::Vector{m});
 
@@ -965,7 +965,7 @@
 
     auto* m = Member(Source{{34, 56}}, "m", ty.mat3x2<f32>(),
                      utils::Vector{
-                         Location(Source{{12, 34}}, 0u),
+                         Location(Source{{12, 34}}, 0_u),
                      });
     Structure("S", utils::Vector{m});
 
@@ -988,7 +988,7 @@
     // }
     auto* output = Structure(
         "Output", utils::Vector{
-                      Member("a", ty.f32(), utils::Vector{Location(0)}),
+                      Member("a", ty.f32(), utils::Vector{Location(0_a)}),
                       Member("b", ty.f32(), utils::Vector{Builtin(ast::BuiltinValue::kFragDepth)}),
                   });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
@@ -1021,7 +1021,7 @@
              Stage(ast::PipelineStage::kVertex),
          },
          utils::Vector{
-             Location(Source{{13, 43}}, 0),
+             Location(Source{{13, 43}}, 0_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -1041,10 +1041,10 @@
     // };
     // @fragment
     // fn main() -> Output { return Output(); }
-    auto* inner =
-        Structure("Inner", utils::Vector{
-                               Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0)}),
-                           });
+    auto* inner = Structure(
+        "Inner", utils::Vector{
+                     Member(Source{{13, 43}}, "a", ty.f32(), utils::Vector{Location(0_a)}),
+                 });
     auto* output = Structure("Output", utils::Vector{
                                            Member(Source{{14, 52}}, "a", ty.Of(inner)),
                                        });
@@ -1072,7 +1072,7 @@
     // }
     auto* output = Structure("Output", utils::Vector{
                                            Member(Source{{13, 43}}, "a", ty.array<f32>(),
-                                                  utils::Vector{Location(Source{{12, 34}}, 0)}),
+                                                  utils::Vector{Location(Source{{12, 34}}, 0_a)}),
                                        });
     Func(Source{{12, 34}}, "main", utils::Empty, ty.Of(output),
          utils::Vector{
@@ -1100,7 +1100,7 @@
              create<ast::WorkgroupAttribute>(Source{{12, 34}}, Expr(1_i)),
          },
          utils::Vector{
-             Location(Source{{12, 34}}, 1),
+             Location(Source{{12, 34}}, 1_a),
          });
 
     EXPECT_FALSE(r()->Resolve());
@@ -1110,7 +1110,7 @@
 TEST_F(LocationAttributeTests, ComputeShaderLocation_Output) {
     auto* input = Param("input", ty.i32(),
                         utils::Vector{
-                            Location(Source{{12, 34}}, 0u),
+                            Location(Source{{12, 34}}, 0_u),
                         });
     Func("main", utils::Vector{input}, ty.void_(), utils::Empty,
          utils::Vector{
@@ -1125,7 +1125,7 @@
 TEST_F(LocationAttributeTests, ComputeShaderLocationStructMember_Output) {
     auto* m = Member("m", ty.i32(),
                      utils::Vector{
-                         Location(Source{{12, 34}}, 0u),
+                         Location(Source{{12, 34}}, 0_u),
                      });
     auto* s = Structure("S", utils::Vector{m});
     Func(Source{{56, 78}}, "main", utils::Empty, ty.Of(s),
@@ -1146,7 +1146,7 @@
 TEST_F(LocationAttributeTests, ComputeShaderLocationStructMember_Input) {
     auto* m = Member("m", ty.i32(),
                      utils::Vector{
-                         Location(Source{{12, 34}}, 0u),
+                         Location(Source{{12, 34}}, 0_u),
                      });
     auto* s = Structure("S", utils::Vector{m});
     auto* input = Param("input", ty.Of(s));
@@ -1168,11 +1168,11 @@
     //         @location(1) param_b : f32) {}
     auto* param_a = Param("param_a", ty.f32(),
                           utils::Vector{
-                              Location(1),
+                              Location(1_a),
                           });
     auto* param_b = Param("param_b", ty.f32(),
                           utils::Vector{
-                              Location(Source{{12, 34}}, 1),
+                              Location(Source{{12, 34}}, 1_a),
                           });
     Func(Source{{12, 34}}, "main",
          utils::Vector{
@@ -1198,12 +1198,12 @@
     // @fragment
     // fn main(param_a : InputA, param_b : InputB) {}
     auto* input_a = Structure("InputA", utils::Vector{
-                                            Member("a", ty.f32(), utils::Vector{Location(1)}),
+                                            Member("a", ty.f32(), utils::Vector{Location(1_a)}),
                                         });
-    auto* input_b =
-        Structure("InputB", utils::Vector{
-                                Member("a", ty.f32(), utils::Vector{Location(Source{{34, 56}}, 1)}),
-                            });
+    auto* input_b = Structure(
+        "InputB", utils::Vector{
+                      Member("a", ty.f32(), utils::Vector{Location(Source{{34, 56}}, 1_a)}),
+                  });
     auto* param_a = Param("param_a", ty.Of(input_a));
     auto* param_b = Param("param_b", ty.Of(input_b));
     Func(Source{{12, 34}}, "main",
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index 675aa41..b1ac41e 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -12584,12 +12584,12 @@
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[635],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpGreaterThanEqual,
   },
   {
     /* [369] */
@@ -12601,55 +12601,55 @@
     /* parameters */ &kParameters[631],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpGreaterThanEqual,
   },
   {
     /* [370] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[641],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpLessThanEqual,
   },
   {
     /* [371] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[639],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpLessThanEqual,
   },
   {
     /* [372] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[649],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpGreaterThan,
   },
   {
     /* [373] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[645],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpGreaterThan,
   },
   {
     /* [374] */
@@ -12824,24 +12824,24 @@
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[657],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpLessThan,
   },
   {
     /* [389] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[15],
+    /* template types */ &kTemplateTypes[13],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[653],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpLessThan,
   },
   {
     /* [390] */
@@ -12992,48 +12992,48 @@
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[16],
+    /* template types */ &kTemplateTypes[18],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[659],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpNotEqual,
   },
   {
     /* [403] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[16],
+    /* template types */ &kTemplateTypes[18],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[599],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpNotEqual,
   },
   {
     /* [404] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[16],
+    /* template types */ &kTemplateTypes[18],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[663],
     /* return matcher indices */ &kMatcherIndices[16],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpEqual,
   },
   {
     /* [405] */
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[16],
+    /* template types */ &kTemplateTypes[18],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[661],
     /* return matcher indices */ &kMatcherIndices[39],
     /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::OpEqual,
   },
   {
     /* [406] */
@@ -14699,42 +14699,42 @@
   },
   {
     /* [10] */
-    /* op ==<T : scalar>(T, T) -> bool */
-    /* op ==<T : scalar, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
+    /* op ==<T : abstract_or_scalar>(T, T) -> bool */
+    /* op ==<T : abstract_or_scalar, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[404],
   },
   {
     /* [11] */
-    /* op !=<T : scalar>(T, T) -> bool */
-    /* op !=<T : scalar, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
+    /* op !=<T : abstract_or_scalar>(T, T) -> bool */
+    /* op !=<T : abstract_or_scalar, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[402],
   },
   {
     /* [12] */
-    /* op <<T : fiu32_f16>(T, T) -> bool */
-    /* op <<T : fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
+    /* op <<T : fia_fiu32_f16>(T, T) -> bool */
+    /* op <<T : fia_fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[388],
   },
   {
     /* [13] */
-    /* op ><T : fiu32_f16>(T, T) -> bool */
-    /* op ><T : fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
+    /* op ><T : fia_fiu32_f16>(T, T) -> bool */
+    /* op ><T : fia_fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[372],
   },
   {
     /* [14] */
-    /* op <=<T : fiu32_f16>(T, T) -> bool */
-    /* op <=<T : fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
+    /* op <=<T : fia_fiu32_f16>(T, T) -> bool */
+    /* op <=<T : fia_fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[370],
   },
   {
     /* [15] */
-    /* op >=<T : fiu32_f16>(T, T) -> bool */
+    /* op >=<T : fia_fiu32_f16>(T, T) -> bool */
     /* op >=<T : fiu32_f16, N : num>(vec<N, T>, vec<N, T>) -> vec<N, bool> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[368],
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 1eb3b35..27dc219 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -382,7 +382,8 @@
     if (is_global) {
         sem = builder_->create<sem::GlobalVariable>(
             v, ty, sem::EvaluationStage::kRuntime, ast::StorageClass::kNone,
-            ast::Access::kUndefined, /* constant_value */ nullptr, sem::BindingPoint{});
+            ast::Access::kUndefined, /* constant_value */ nullptr, sem::BindingPoint{},
+            std::nullopt);
     } else {
         sem = builder_->create<sem::LocalVariable>(v, ty, sem::EvaluationStage::kRuntime,
                                                    ast::StorageClass::kNone,
@@ -437,11 +438,11 @@
 
     auto* sem = builder_->create<sem::GlobalVariable>(
         v, ty, sem::EvaluationStage::kOverride, ast::StorageClass::kNone, ast::Access::kUndefined,
-        /* constant_value */ nullptr, sem::BindingPoint{});
+        /* constant_value */ nullptr, sem::BindingPoint{}, std::nullopt);
     sem->SetConstructor(rhs);
 
     if (auto* id_attr = ast::GetAttribute<ast::IdAttribute>(v->attributes)) {
-        auto* materialize = Materialize(Expression(id_attr->value));
+        auto* materialize = Materialize(Expression(id_attr->expr));
         if (!materialize) {
             return nullptr;
         }
@@ -521,7 +522,7 @@
 
     auto* sem = is_global ? static_cast<sem::Variable*>(builder_->create<sem::GlobalVariable>(
                                 c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
-                                ast::Access::kUndefined, value, sem::BindingPoint{}))
+                                ast::Access::kUndefined, value, sem::BindingPoint{}, std::nullopt))
                           : static_cast<sem::Variable*>(builder_->create<sem::LocalVariable>(
                                 c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
                                 ast::Access::kUndefined, current_statement_, value));
@@ -606,7 +607,7 @@
             uint32_t binding = 0;
             {
                 auto* attr = ast::GetAttribute<ast::BindingAttribute>(var->attributes);
-                auto* materialize = Materialize(Expression(attr->value));
+                auto* materialize = Materialize(Expression(attr->expr));
                 if (!materialize) {
                     return nullptr;
                 }
@@ -622,7 +623,7 @@
             uint32_t group = 0;
             {
                 auto* attr = ast::GetAttribute<ast::GroupAttribute>(var->attributes);
-                auto* materialize = Materialize(Expression(attr->value));
+                auto* materialize = Materialize(Expression(attr->expr));
                 if (!materialize) {
                     return nullptr;
                 }
@@ -636,9 +637,25 @@
             }
             binding_point = {group, binding};
         }
-        sem = builder_->create<sem::GlobalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
-                                                    storage_class, access,
-                                                    /* constant_value */ nullptr, binding_point);
+
+        std::optional<uint32_t> location;
+        if (auto* attr = ast::GetAttribute<ast::LocationAttribute>(var->attributes)) {
+            auto* materialize = Materialize(Expression(attr->expr));
+            if (!materialize) {
+                return nullptr;
+            }
+            auto* c = materialize->ConstantValue();
+            if (!c) {
+                // TODO(crbug.com/tint/1633): Add error message about invalid materialization
+                // when location can be an expression.
+                return nullptr;
+            }
+            location = c->As<uint32_t>();
+        }
+
+        sem = builder_->create<sem::GlobalVariable>(
+            var, var_ty, sem::EvaluationStage::kRuntime, storage_class, access,
+            /* constant_value */ nullptr, binding_point, location);
 
     } else {
         sem = builder_->create<sem::LocalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
@@ -688,7 +705,7 @@
     if (param->HasBindingPoint()) {
         {
             auto* attr = ast::GetAttribute<ast::BindingAttribute>(param->attributes);
-            auto* materialize = Materialize(Expression(attr->value));
+            auto* materialize = Materialize(Expression(attr->expr));
             if (!materialize) {
                 return nullptr;
             }
@@ -702,7 +719,7 @@
         }
         {
             auto* attr = ast::GetAttribute<ast::GroupAttribute>(param->attributes);
-            auto* materialize = Materialize(Expression(attr->value));
+            auto* materialize = Materialize(Expression(attr->expr));
             if (!materialize) {
                 return nullptr;
             }
@@ -718,7 +735,17 @@
 
     std::optional<uint32_t> location;
     if (auto* l = ast::GetAttribute<ast::LocationAttribute>(param->attributes)) {
-        location = l->value;
+        auto* materialize = Materialize(Expression(l->expr));
+        if (!materialize) {
+            return nullptr;
+        }
+        auto* c = materialize->ConstantValue();
+        if (!c) {
+            // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
+            // location can be an expression.
+            return nullptr;
+        }
+        location = c->As<uint32_t>();
     }
 
     auto* sem = builder_->create<sem::Parameter>(
@@ -917,7 +944,17 @@
         Mark(attr);
 
         if (auto* a = attr->As<ast::LocationAttribute>()) {
-            return_location = a->value;
+            auto* materialize = Materialize(Expression(a->expr));
+            if (!materialize) {
+                return nullptr;
+            }
+            auto* c = materialize->ConstantValue();
+            if (!c) {
+                // TODO(crbug.com/tint/1633): Add error message about invalid materialization when
+                // location can be an expression.
+                return nullptr;
+            }
+            return_location = c->As<uint32_t>();
         }
     }
     if (!validator_.NoDuplicateAttributes(decl->attributes)) {
@@ -2510,6 +2547,7 @@
             if (!op.result) {
                 return nullptr;
             }
+            ty = op.result;
             if (ShouldMaterializeArgument(op.parameter)) {
                 expr = Materialize(expr, op.parameter);
                 if (!expr) {
@@ -2530,7 +2568,6 @@
                     stage = sem::EvaluationStage::kRuntime;
                 }
             }
-            ty = op.result;
             break;
         }
     }
@@ -2774,13 +2811,13 @@
                 align = 1;
                 has_offset_attr = true;
             } else if (auto* a = attr->As<ast::StructMemberAlignAttribute>()) {
-                auto* materialized = Materialize(Expression(a->align));
+                auto* materialized = Materialize(Expression(a->expr));
                 if (!materialized) {
                     return nullptr;
                 }
                 auto const_value = materialized->ConstantValue();
                 if (!const_value) {
-                    AddError("'align' must be constant expression", a->align->source);
+                    AddError("'align' must be constant expression", a->expr->source);
                     return nullptr;
                 }
                 auto value = const_value->As<AInt>();
@@ -2801,7 +2838,17 @@
                 size = s->size;
                 has_size_attr = true;
             } else if (auto* l = attr->As<ast::LocationAttribute>()) {
-                location = l->value;
+                auto* materialize = Materialize(Expression(l->expr));
+                if (!materialize) {
+                    return nullptr;
+                }
+                auto* c = materialize->ConstantValue();
+                if (!c) {
+                    // TODO(crbug.com/tint/1633): Add error message about invalid materialization
+                    // when location can be an expression.
+                    return nullptr;
+                }
+                location = c->As<uint32_t>();
             }
         }
 
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index 523dd89..1a3c623 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -774,9 +774,9 @@
 }
 
 TEST_F(ResolverTest, Function_Parameters_Locations) {
-    auto* param_a = Param("a", ty.f32(), utils::Vector{Location(3)});
+    auto* param_a = Param("a", ty.f32(), utils::Vector{Location(3_a)});
     auto* param_b = Param("b", ty.u32(), utils::Vector{Builtin(ast::BuiltinValue::kVertexIndex)});
-    auto* param_c = Param("c", ty.u32(), utils::Vector{Location(1)});
+    auto* param_c = Param("c", ty.u32(), utils::Vector{Location(1_a)});
 
     GlobalVar("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
     auto* func = Func("my_func",
@@ -806,6 +806,18 @@
     EXPECT_EQ(1u, func_sem->Parameters()[2]->Location());
 }
 
+TEST_F(ResolverTest, Function_GlobalVariable_Location) {
+    auto* var = GlobalVar(
+        "my_vec", ty.vec4<f32>(), ast::StorageClass::kIn,
+        utils::Vector{Location(3_a), Disable(ast::DisabledValidation::kIgnoreStorageClass)});
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get<sem::GlobalVariable>(var);
+    ASSERT_NE(sem, nullptr);
+    EXPECT_EQ(3u, sem->Location());
+}
+
 TEST_F(ResolverTest, Function_RegisterInputOutputVariables) {
     auto* s = Structure("S", utils::Vector{Member("m", ty.u32())});
 
@@ -844,7 +856,7 @@
                           Stage(ast::PipelineStage::kFragment),
                       },
                       utils::Vector{
-                          Location(2),
+                          Location(2_a),
                       });
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();
diff --git a/src/tint/resolver/resolver_test_helper.h b/src/tint/resolver/resolver_test_helper.h
index 6641176..77d1481 100644
--- a/src/tint/resolver/resolver_test_helper.h
+++ b/src/tint/resolver/resolver_test_helper.h
@@ -751,6 +751,12 @@
 template <typename T>
 constexpr bool IsValue = detail::IsValue<T>::value;
 
+/// Returns the friendly name of ValueT
+template <typename ValueT, typename = traits::EnableIf<IsValue<ValueT>>>
+const char* FriendlyName() {
+    return tint::FriendlyName<typename ValueT::ElementType>();
+}
+
 /// Creates a `Value<T>` from a scalar `v`
 template <typename T>
 auto Val(T v) {
diff --git a/src/tint/resolver/struct_pipeline_stage_use_test.cc b/src/tint/resolver/struct_pipeline_stage_use_test.cc
index c8e77ea..107b241 100644
--- a/src/tint/resolver/struct_pipeline_stage_use_test.cc
+++ b/src/tint/resolver/struct_pipeline_stage_use_test.cc
@@ -29,7 +29,7 @@
 using ResolverPipelineStageUseTest = ResolverTest;
 
 TEST_F(ResolverPipelineStageUseTest, UnusedStruct) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     ASSERT_TRUE(r()->Resolve()) << r()->error();
 
@@ -39,7 +39,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsNonEntryPointParam) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     Func("foo", utils::Vector{Param("param", ty.Of(s))}, ty.void_(), utils::Empty, utils::Empty);
 
@@ -51,7 +51,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsNonEntryPointReturnType) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     Func("foo", utils::Empty, ty.Of(s), utils::Vector{Return(Construct(ty.Of(s), Expr(0_f)))},
          utils::Empty);
@@ -64,7 +64,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsVertexShaderParam) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     Func("main", utils::Vector{Param("param", ty.Of(s))}, ty.vec4<f32>(),
          utils::Vector{Return(Construct(ty.vec4<f32>()))},
@@ -96,7 +96,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsFragmentShaderParam) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     Func("main", utils::Vector{Param("param", ty.Of(s))}, ty.void_(), utils::Empty,
          utils::Vector{Stage(ast::PipelineStage::kFragment)});
@@ -110,7 +110,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsFragmentShaderReturnType) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
 
     Func("main", utils::Empty, ty.Of(s), utils::Vector{Return(Construct(ty.Of(s), Expr(0_f)))},
          utils::Vector{Stage(ast::PipelineStage::kFragment)});
@@ -160,7 +160,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderParamViaAlias) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
     auto* s_alias = Alias("S_alias", ty.Of(s));
 
     Func("main", utils::Vector{Param("param", ty.Of(s_alias))}, ty.void_(), utils::Empty,
@@ -175,7 +175,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderParamLocationSet) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3_a)})});
 
     Func("main", utils::Vector{Param("param", ty.Of(s))}, ty.void_(), utils::Empty,
          utils::Vector{Stage(ast::PipelineStage::kFragment)});
@@ -189,7 +189,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderReturnTypeViaAlias) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0_a)})});
     auto* s_alias = Alias("S_alias", ty.Of(s));
 
     Func("main", utils::Empty, ty.Of(s_alias),
@@ -205,7 +205,7 @@
 }
 
 TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderReturnTypeLocationSet) {
-    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3)})});
+    auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3_a)})});
 
     Func("main", utils::Empty, ty.Of(s), utils::Vector{Return(Construct(ty.Of(s), Expr(0_f)))},
          utils::Vector{Stage(ast::PipelineStage::kFragment)});
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index caca32e..1812789 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -121,12 +121,13 @@
 }
 
 // Helper to stringify a pipeline IO attribute.
-std::string attr_to_str(const ast::Attribute* attr) {
+std::string attr_to_str(const ast::Attribute* attr,
+                        std::optional<uint32_t> location = std::nullopt) {
     std::stringstream str;
     if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
         str << "builtin(" << builtin->builtin << ")";
-    } else if (auto* location = attr->As<ast::LocationAttribute>()) {
-        str << "location(" << location->value << ")";
+    } else if (attr->Is<ast::LocationAttribute>()) {
+        str << "location(" << location.value() << ")";
     }
     return str.str();
 }
@@ -601,7 +602,8 @@
                 if (!attr->IsAnyOf<ast::BindingAttribute, ast::GroupAttribute,
                                    ast::InternalAttribute>() &&
                     (!is_shader_io_attribute || !has_io_storage_class)) {
-                    AddError("attribute is not valid for module-scope 'var'", attr->source);
+                    AddError("attribute '" + attr->Name() + "' is not valid for module-scope 'var'",
+                             attr->source);
                     return false;
                 }
             }
@@ -1122,7 +1124,8 @@
     auto validate_entry_point_attributes_inner = [&](utils::VectorRef<const ast::Attribute*> attrs,
                                                      const sem::Type* ty, Source source,
                                                      ParamOrRetType param_or_ret,
-                                                     bool is_struct_member) {
+                                                     bool is_struct_member,
+                                                     std::optional<uint32_t> location) {
         // Temporally forbid using f16 types in entry point IO.
         // TODO(tint:1473, tint:1502): Remove this error after f16 is supported in entry point
         // IO.
@@ -1142,7 +1145,7 @@
             if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
                 if (pipeline_io_attribute) {
                     AddError("multiple entry point IO attributes", attr->source);
-                    AddNote("previously consumed " + attr_to_str(pipeline_io_attribute),
+                    AddNote("previously consumed " + attr_to_str(pipeline_io_attribute, location),
                             pipeline_io_attribute->source);
                     return false;
                 }
@@ -1161,7 +1164,7 @@
                     return false;
                 }
                 builtins.emplace(builtin->builtin);
-            } else if (auto* location = attr->As<ast::LocationAttribute>()) {
+            } else if (auto* loc_attr = attr->As<ast::LocationAttribute>()) {
                 if (pipeline_io_attribute) {
                     AddError("multiple entry point IO attributes", attr->source);
                     AddNote("previously consumed " + attr_to_str(pipeline_io_attribute),
@@ -1172,7 +1175,13 @@
 
                 bool is_input = param_or_ret == ParamOrRetType::kParameter;
 
-                if (!LocationAttribute(location, ty, locations, stage, source, is_input)) {
+                if (!location.has_value()) {
+                    TINT_ICE(Resolver, diagnostics_) << "Location has no value";
+                    return false;
+                }
+
+                if (!LocationAttribute(loc_attr, location.value(), ty, locations, stage, source,
+                                       is_input)) {
                     return false;
                 }
             } else if (auto* interpolate = attr->As<ast::InterpolateAttribute>()) {
@@ -1265,9 +1274,10 @@
     // Outer lambda for validating the entry point attributes for a type.
     auto validate_entry_point_attributes = [&](utils::VectorRef<const ast::Attribute*> attrs,
                                                const sem::Type* ty, Source source,
-                                               ParamOrRetType param_or_ret) {
+                                               ParamOrRetType param_or_ret,
+                                               std::optional<uint32_t> location) {
         if (!validate_entry_point_attributes_inner(attrs, ty, source, param_or_ret,
-                                                   /*is_struct_member*/ false)) {
+                                                   /*is_struct_member*/ false, location)) {
             return false;
         }
 
@@ -1276,7 +1286,7 @@
                 if (!validate_entry_point_attributes_inner(
                         member->Declaration()->attributes, member->Type(),
                         member->Declaration()->source, param_or_ret,
-                        /*is_struct_member*/ true)) {
+                        /*is_struct_member*/ true, member->Location())) {
                     AddNote("while analysing entry point '" + symbols_.NameFor(decl->symbol) + "'",
                             decl->source);
                     return false;
@@ -1290,7 +1300,8 @@
     for (auto* param : func->Parameters()) {
         auto* param_decl = param->Declaration();
         if (!validate_entry_point_attributes(param_decl->attributes, param->Type(),
-                                             param_decl->source, ParamOrRetType::kParameter)) {
+                                             param_decl->source, ParamOrRetType::kParameter,
+                                             param->Location())) {
             return false;
         }
     }
@@ -1303,7 +1314,8 @@
 
     if (!func->ReturnType()->Is<sem::Void>()) {
         if (!validate_entry_point_attributes(decl->return_type_attributes, func->ReturnType(),
-                                             decl->source, ParamOrRetType::kReturnType)) {
+                                             decl->source, ParamOrRetType::kReturnType,
+                                             func->ReturnLocation())) {
             return false;
         }
     }
@@ -2176,8 +2188,9 @@
                 invariant_attribute = invariant;
             } else if (auto* location = attr->As<ast::LocationAttribute>()) {
                 has_location = true;
-                if (!LocationAttribute(location, member->Type(), locations, stage,
-                                       member->Declaration()->source)) {
+                TINT_ASSERT(Resolver, member->Location().has_value());
+                if (!LocationAttribute(location, member->Location().value(), member->Type(),
+                                       locations, stage, member->Declaration()->source)) {
                     return false;
                 }
             } else if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
@@ -2219,7 +2232,8 @@
     return true;
 }
 
-bool Validator::LocationAttribute(const ast::LocationAttribute* location,
+bool Validator::LocationAttribute(const ast::LocationAttribute* loc_attr,
+                                  uint32_t location,
                                   const sem::Type* type,
                                   std::unordered_set<uint32_t>& locations,
                                   ast::PipelineStage stage,
@@ -2227,7 +2241,7 @@
                                   const bool is_input) const {
     std::string inputs_or_output = is_input ? "inputs" : "output";
     if (stage == ast::PipelineStage::kCompute) {
-        AddError("attribute is not valid for compute shader " + inputs_or_output, location->source);
+        AddError("attribute is not valid for compute shader " + inputs_or_output, loc_attr->source);
         return false;
     }
 
@@ -2238,15 +2252,16 @@
         AddNote(
             "'location' attribute must only be applied to declarations of "
             "numeric scalar or numeric vector type",
-            location->source);
+            loc_attr->source);
         return false;
     }
 
-    if (locations.count(location->value)) {
-        AddError(attr_to_str(location) + " attribute appears multiple times", location->source);
+    if (locations.count(location)) {
+        AddError(attr_to_str(loc_attr, location) + " attribute appears multiple times",
+                 loc_attr->source);
         return false;
     }
-    locations.emplace(location->value);
+    locations.emplace(location);
 
     return true;
 }
diff --git a/src/tint/resolver/validator.h b/src/tint/resolver/validator.h
index 8bec86f..a00f6ab 100644
--- a/src/tint/resolver/validator.h
+++ b/src/tint/resolver/validator.h
@@ -273,14 +273,16 @@
     bool LocalVariable(const sem::Variable* v) const;
 
     /// Validates a location attribute
-    /// @param location the location attribute to validate
+    /// @param loc_attr the location attribute to validate
+    /// @param location the location value
     /// @param type the variable type
     /// @param locations the set of locations in the module
     /// @param stage the current pipeline stage
     /// @param source the source of the attribute
     /// @param is_input true if this is an input variable
     /// @returns true on success, false otherwise.
-    bool LocationAttribute(const ast::LocationAttribute* location,
+    bool LocationAttribute(const ast::LocationAttribute* loc_attr,
+                           uint32_t location,
                            const sem::Type* type,
                            std::unordered_set<uint32_t>& locations,
                            ast::PipelineStage stage,
diff --git a/src/tint/sem/sem_struct_test.cc b/src/tint/sem/sem_struct_test.cc
index 453b8f7..9746589 100644
--- a/src/tint/sem/sem_struct_test.cc
+++ b/src/tint/sem/sem_struct_test.cc
@@ -19,6 +19,7 @@
 namespace tint::sem {
 namespace {
 
+using namespace tint::number_suffixes;  // NOLINT
 using StructTest = TestHelper;
 
 TEST_F(StructTest, Creation) {
@@ -107,7 +108,7 @@
 
 TEST_F(StructTest, Location) {
     auto* st = Structure("st", utils::Vector{
-                                   Member("a", ty.i32(), utils::Vector{Location(1u)}),
+                                   Member("a", ty.i32(), utils::Vector{Location(1_u)}),
                                    Member("b", ty.u32()),
                                });
 
diff --git a/src/tint/sem/variable.cc b/src/tint/sem/variable.cc
index 67e7945..6dcec63 100644
--- a/src/tint/sem/variable.cc
+++ b/src/tint/sem/variable.cc
@@ -61,9 +61,11 @@
                                ast::StorageClass storage_class,
                                ast::Access access,
                                const Constant* constant_value,
-                               sem::BindingPoint binding_point)
+                               sem::BindingPoint binding_point,
+                               std::optional<uint32_t> location)
     : Base(declaration, type, stage, storage_class, access, constant_value),
-      binding_point_(binding_point) {}
+      binding_point_(binding_point),
+      location_(location) {}
 
 GlobalVariable::~GlobalVariable() = default;
 
diff --git a/src/tint/sem/variable.h b/src/tint/sem/variable.h
index a0eb75e..5ea70b9 100644
--- a/src/tint/sem/variable.h
+++ b/src/tint/sem/variable.h
@@ -153,13 +153,18 @@
     /// @param access the variable access control type
     /// @param constant_value the constant value for the variable. May be null
     /// @param binding_point the optional resource binding point of the variable
+    /// @param location the location value if provided
+    ///
+    /// Note, a GlobalVariable generally doesn't have a `location` in WGSL, as it isn't allowed by
+    /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
     GlobalVariable(const ast::Variable* declaration,
                    const sem::Type* type,
                    EvaluationStage stage,
                    ast::StorageClass storage_class,
                    ast::Access access,
                    const Constant* constant_value,
-                   sem::BindingPoint binding_point = {});
+                   sem::BindingPoint binding_point = {},
+                   std::optional<uint32_t> location = std::nullopt);
 
     /// Destructor
     ~GlobalVariable() override;
@@ -173,10 +178,14 @@
     /// @returns the pipeline constant ID associated with the variable
     tint::OverrideId OverrideId() const { return override_id_; }
 
+    /// @returns the location value for the parameter, if set
+    std::optional<uint32_t> Location() const { return location_; }
+
   private:
     const sem::BindingPoint binding_point_;
 
     tint::OverrideId override_id_;
+    std::optional<uint32_t> location_;
 };
 
 /// Parameter is a function parameter
diff --git a/src/tint/transform/canonicalize_entry_point_io.cc b/src/tint/transform/canonicalize_entry_point_io.cc
index 31feb86..b08d44d 100644
--- a/src/tint/transform/canonicalize_entry_point_io.cc
+++ b/src/tint/transform/canonicalize_entry_point_io.cc
@@ -37,21 +37,32 @@
 
 namespace {
 
-// Comparison function used to reorder struct members such that all members with
-// location attributes appear first (ordered by location slot), followed by
-// those with builtin attributes.
-bool StructMemberComparator(const ast::StructMember* a, const ast::StructMember* b) {
-    auto* a_loc = ast::GetAttribute<ast::LocationAttribute>(a->attributes);
-    auto* b_loc = ast::GetAttribute<ast::LocationAttribute>(b->attributes);
-    auto* a_blt = ast::GetAttribute<ast::BuiltinAttribute>(a->attributes);
-    auto* b_blt = ast::GetAttribute<ast::BuiltinAttribute>(b->attributes);
+/// Info for a struct member
+struct MemberInfo {
+    /// The struct member item
+    const ast::StructMember* member;
+    /// The struct member location if provided
+    std::optional<uint32_t> location;
+};
+
+/// Comparison function used to reorder struct members such that all members with
+/// location attributes appear first (ordered by location slot), followed by
+/// those with builtin attributes.
+/// @param a a struct member
+/// @param b another struct member
+/// @returns true if a comes before b
+bool StructMemberComparator(const MemberInfo& a, const MemberInfo& b) {
+    auto* a_loc = ast::GetAttribute<ast::LocationAttribute>(a.member->attributes);
+    auto* b_loc = ast::GetAttribute<ast::LocationAttribute>(b.member->attributes);
+    auto* a_blt = ast::GetAttribute<ast::BuiltinAttribute>(a.member->attributes);
+    auto* b_blt = ast::GetAttribute<ast::BuiltinAttribute>(b.member->attributes);
     if (a_loc) {
         if (!b_loc) {
             // `a` has location attribute and `b` does not: `a` goes first.
             return true;
         }
         // Both have location attributes: smallest goes first.
-        return a_loc->value < b_loc->value;
+        return a.location < b.location;
     } else {
         if (b_loc) {
             // `b` has location attribute and `a` does not: `b` goes first.
@@ -88,6 +99,8 @@
         utils::Vector<const ast::Attribute*, 2> attributes;
         /// The value itself.
         const ast::Expression* value;
+        /// The output location.
+        std::optional<uint32_t> location;
     };
 
     /// The clone context.
@@ -101,14 +114,15 @@
 
     /// The new entry point wrapper function's parameters.
     utils::Vector<const ast::Parameter*, 8> wrapper_ep_parameters;
+
     /// The members of the wrapper function's struct parameter.
-    utils::Vector<const ast::StructMember*, 8> wrapper_struct_param_members;
+    utils::Vector<MemberInfo, 8> wrapper_struct_param_members;
     /// The name of the wrapper function's struct parameter.
     Symbol wrapper_struct_param_name;
     /// The parameters that will be passed to the original function.
     utils::Vector<const ast::Expression*, 8> inner_call_parameters;
     /// The members of the wrapper function's struct return type.
-    utils::Vector<const ast::StructMember*, 8> wrapper_struct_output_members;
+    utils::Vector<MemberInfo, 8> wrapper_struct_output_members;
     /// The wrapper function output values.
     utils::Vector<OutputValue, 8> wrapper_output_values;
     /// The body of the wrapper function.
@@ -153,10 +167,12 @@
     /// Add a shader input to the entry point.
     /// @param name the name of the shader input
     /// @param type the type of the shader input
+    /// @param location the location if provided
     /// @param attributes the attributes to apply to the shader input
     /// @returns an expression which evaluates to the value of the shader input
     const ast::Expression* AddInput(std::string name,
                                     const sem::Type* type,
+                                    std::optional<uint32_t> location,
                                     utils::Vector<const ast::Attribute*, 8> attributes) {
         auto* ast_type = CreateASTTypeFor(ctx, type);
         if (cfg.shader_style == ShaderStyle::kSpirv || cfg.shader_style == ShaderStyle::kGlsl) {
@@ -214,7 +230,7 @@
             Symbol symbol = input_names.emplace(name).second ? ctx.dst->Symbols().Register(name)
                                                              : ctx.dst->Symbols().New(name);
             wrapper_struct_param_members.Push(
-                ctx.dst->Member(symbol, ast_type, std::move(attributes)));
+                {ctx.dst->Member(symbol, ast_type, std::move(attributes)), location});
             return ctx.dst->MemberAccessor(InputStructSymbol(), symbol);
         }
     }
@@ -222,10 +238,12 @@
     /// Add a shader output to the entry point.
     /// @param name the name of the shader output
     /// @param type the type of the shader output
+    /// @param location the location if provided
     /// @param attributes the attributes to apply to the shader output
     /// @param value the value of the shader output
     void AddOutput(std::string name,
                    const sem::Type* type,
+                   std::optional<uint32_t> location,
                    utils::Vector<const ast::Attribute*, 8> attributes,
                    const ast::Expression* value) {
         // Vulkan requires that integer user-defined vertex outputs are always decorated with
@@ -256,6 +274,7 @@
         output.type = CreateASTTypeFor(ctx, type);
         output.attributes = std::move(attributes);
         output.value = value;
+        output.location = location;
         wrapper_output_values.Push(output);
     }
 
@@ -280,7 +299,7 @@
         }
 
         auto name = ctx.src->Symbols().NameFor(param->Declaration()->symbol);
-        auto* input_expr = AddInput(name, param->Type(), std::move(attributes));
+        auto* input_expr = AddInput(name, param->Type(), param->Location(), std::move(attributes));
         inner_call_parameters.Push(input_expr);
     }
 
@@ -308,7 +327,8 @@
             auto name = ctx.src->Symbols().NameFor(member_ast->symbol);
 
             auto attributes = CloneShaderIOAttributes(member_ast->attributes, do_interpolate);
-            auto* input_expr = AddInput(name, member->Type(), std::move(attributes));
+            auto* input_expr =
+                AddInput(name, member->Type(), member->Location(), std::move(attributes));
             inner_struct_values.Push(input_expr);
         }
 
@@ -337,7 +357,7 @@
                 auto attributes = CloneShaderIOAttributes(member_ast->attributes, do_interpolate);
 
                 // Extract the original structure member.
-                AddOutput(name, member->Type(), std::move(attributes),
+                AddOutput(name, member->Type(), member->Location(), std::move(attributes),
                           ctx.dst->MemberAccessor(original_result, name));
             }
         } else if (!inner_ret_type->Is<sem::Void>()) {
@@ -345,8 +365,8 @@
                 CloneShaderIOAttributes(func_ast->return_type_attributes, do_interpolate);
 
             // Propagate the non-struct return value as is.
-            AddOutput("value", func_sem->ReturnType(), std::move(attributes),
-                      ctx.dst->Expr(original_result));
+            AddOutput("value", func_sem->ReturnType(), func_sem->ReturnLocation(),
+                      std::move(attributes), ctx.dst->Expr(original_result));
         }
     }
 
@@ -365,7 +385,7 @@
 
         // No existing sample mask builtin was found, so create a new output value
         // using the fixed sample mask.
-        AddOutput("fixed_sample_mask", ctx.dst->create<sem::U32>(),
+        AddOutput("fixed_sample_mask", ctx.dst->create<sem::U32>(), std::nullopt,
                   {ctx.dst->Builtin(ast::BuiltinValue::kSampleMask)},
                   ctx.dst->Expr(u32(cfg.fixed_sample_mask)));
     }
@@ -373,7 +393,7 @@
     /// Add a point size builtin to the wrapper function output.
     void AddVertexPointSize() {
         // Create a new output value and assign it a literal 1.0 value.
-        AddOutput("vertex_point_size", ctx.dst->create<sem::F32>(),
+        AddOutput("vertex_point_size", ctx.dst->create<sem::F32>(), std::nullopt,
                   {ctx.dst->Builtin(ast::BuiltinValue::kPointSize)}, ctx.dst->Expr(1_f));
     }
 
@@ -392,10 +412,14 @@
         std::sort(wrapper_struct_param_members.begin(), wrapper_struct_param_members.end(),
                   StructMemberComparator);
 
+        utils::Vector<const ast::StructMember*, 8> members;
+        for (auto& mem : wrapper_struct_param_members) {
+            members.Push(mem.member);
+        }
+
         // Create the new struct type.
         auto struct_name = ctx.dst->Sym();
-        auto* in_struct =
-            ctx.dst->create<ast::Struct>(struct_name, wrapper_struct_param_members, utils::Empty);
+        auto* in_struct = ctx.dst->create<ast::Struct>(struct_name, members, utils::Empty);
         ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast, in_struct);
 
         // Create a new function parameter using this struct type.
@@ -423,7 +447,8 @@
             member_names.insert(ctx.dst->Symbols().NameFor(name));
 
             wrapper_struct_output_members.Push(
-                ctx.dst->Member(name, outval.type, std::move(outval.attributes)));
+                {ctx.dst->Member(name, outval.type, std::move(outval.attributes)),
+                 outval.location});
             assignments.Push(
                 ctx.dst->Assign(ctx.dst->MemberAccessor(wrapper_result, name), outval.value));
         }
@@ -432,9 +457,13 @@
         std::sort(wrapper_struct_output_members.begin(), wrapper_struct_output_members.end(),
                   StructMemberComparator);
 
+        utils::Vector<const ast::StructMember*, 8> members;
+        for (auto& mem : wrapper_struct_output_members) {
+            members.Push(mem.member);
+        }
+
         // Create the new struct type.
-        auto* out_struct = ctx.dst->create<ast::Struct>(
-            ctx.dst->Sym(), wrapper_struct_output_members, utils::Empty);
+        auto* out_struct = ctx.dst->create<ast::Struct>(ctx.dst->Sym(), members, utils::Empty);
         ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast, out_struct);
 
         // Create the output struct object, assign its members, and return it.
diff --git a/src/tint/transform/single_entry_point.cc b/src/tint/transform/single_entry_point.cc
index 133c836..8d26a7f 100644
--- a/src/tint/transform/single_entry_point.cc
+++ b/src/tint/transform/single_entry_point.cc
@@ -86,11 +86,17 @@
                     ctx.dst->AST().AddGlobalVariable(ctx.Clone(override));
                 }
             },
-            [&](const ast::Variable* v) {  // var, let
-                if (referenced_vars.count(v)) {
-                    ctx.dst->AST().AddGlobalVariable(ctx.Clone(v));
+            [&](const ast::Var* var) {
+                if (referenced_vars.count(var)) {
+                    ctx.dst->AST().AddGlobalVariable(ctx.Clone(var));
                 }
             },
+            [&](const ast::Const* c) {
+                // Always keep 'const' declarations, as these can be used by attributes and array
+                // sizes, which are not tracked as transitively used by functions. They also don't
+                // typically get emitted by the backend unless they're actually used.
+                ctx.dst->AST().AddGlobalVariable(ctx.Clone(c));
+            },
             [&](const ast::Function* func) {
                 if (sem.Get(func)->HasAncestorEntryPoint(entry_point->symbol)) {
                     ctx.dst->AST().AddFunction(ctx.Clone(func));
diff --git a/src/tint/transform/single_entry_point_test.cc b/src/tint/transform/single_entry_point_test.cc
index 020bd65..7451090 100644
--- a/src/tint/transform/single_entry_point_test.cc
+++ b/src/tint/transform/single_entry_point_test.cc
@@ -217,8 +217,14 @@
 )";
 
     auto* expect = R"(
+const a : f32 = 1.0;
+
+const b : f32 = 1.0;
+
 const c : f32 = 1.0;
 
+const d : f32 = 1.0;
+
 @compute @workgroup_size(1)
 fn comp_main1() {
   let local_c : f32 = c;
@@ -536,5 +542,28 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(SingleEntryPointTest, GlobalConstUsedAsArraySize) {
+    // See crbug.com/tint/1598
+    auto* src = R"(
+const MY_SIZE = 5u;
+
+type Arr = array<i32, MY_SIZE>;
+
+@fragment
+fn main() {
+}
+)";
+
+    auto* expect = src;
+
+    SingleEntryPoint::Config cfg("main");
+
+    DataMap data;
+    data.Add<SingleEntryPoint::Config>(cfg);
+    auto got = Run<SingleEntryPoint>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::transform
diff --git a/src/tint/transform/substitute_override.h b/src/tint/transform/substitute_override.h
index 9ea315d..940e11d 100644
--- a/src/tint/transform/substitute_override.h
+++ b/src/tint/transform/substitute_override.h
@@ -20,6 +20,7 @@
 
 #include "tint/override_id.h"
 
+#include "src/tint/reflection.h"
 #include "src/tint/transform/transform.h"
 
 namespace tint::transform {
@@ -63,6 +64,9 @@
         /// The value is always a double coming into the transform and will be
         /// converted to the correct type through and initializer.
         std::unordered_map<OverrideId, double> map;
+
+        /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+        TINT_REFLECT(map);
     };
 
     /// Constructor
diff --git a/src/tint/transform/vertex_pulling.cc b/src/tint/transform/vertex_pulling.cc
index 40d8d30..3c0dce9 100644
--- a/src/tint/transform/vertex_pulling.cc
+++ b/src/tint/transform/vertex_pulling.cc
@@ -692,7 +692,7 @@
     /// @param func the entry point function
     /// @param param the parameter to process
     void ProcessNonStructParameter(const ast::Function* func, const ast::Parameter* param) {
-        if (auto* location = ast::GetAttribute<ast::LocationAttribute>(param->attributes)) {
+        if (ast::HasAttribute<ast::LocationAttribute>(param->attributes)) {
             // Create a function-scope variable to replace the parameter.
             auto func_var_sym = ctx.Clone(param->symbol);
             auto* func_var_type = ctx.Clone(param->type);
@@ -701,8 +701,15 @@
             // Capture mapping from location to the new variable.
             LocationInfo info;
             info.expr = [this, func_var]() { return ctx.dst->Expr(func_var); };
-            info.type = ctx.src->Sem().Get(param)->Type();
-            location_info[location->value] = info;
+
+            auto* sem = ctx.src->Sem().Get<sem::Parameter>(param);
+            info.type = sem->Type();
+
+            if (!sem->Location().has_value()) {
+                TINT_ICE(Transform, ctx.dst->Diagnostics()) << "Location missing value";
+                return;
+            }
+            location_info[sem->Location().value()] = info;
         } else if (auto* builtin = ast::GetAttribute<ast::BuiltinAttribute>(param->attributes)) {
             // Check for existing vertex_index and instance_index builtins.
             if (builtin->builtin == ast::BuiltinValue::kVertexIndex) {
@@ -742,12 +749,16 @@
                 return ctx.dst->MemberAccessor(param_sym, member_sym);
             };
 
-            if (auto* location = ast::GetAttribute<ast::LocationAttribute>(member->attributes)) {
+            if (ast::HasAttribute<ast::LocationAttribute>(member->attributes)) {
                 // Capture mapping from location to struct member.
                 LocationInfo info;
                 info.expr = member_expr;
-                info.type = ctx.src->Sem().Get(member)->Type();
-                location_info[location->value] = info;
+
+                auto* sem = ctx.src->Sem().Get(member);
+                info.type = sem->Type();
+
+                TINT_ASSERT(Transform, sem->Location().has_value());
+                location_info[sem->Location().value()] = info;
                 has_locations = true;
             } else if (auto* builtin =
                            ast::GetAttribute<ast::BuiltinAttribute>(member->attributes)) {
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 8b2d898..132f542 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -1856,7 +1856,7 @@
     return Switch(
         global,  //
         [&](const ast::Var* var) {
-            auto* sem = builder_.Sem().Get(global);
+            auto* sem = builder_.Sem().Get<sem::GlobalVariable>(global);
             switch (sem->StorageClass()) {
                 case ast::StorageClass::kUniform:
                     return EmitUniformVariable(var, sem);
@@ -2005,7 +2005,7 @@
     return true;
 }
 
-bool GeneratorImpl::EmitIOVariable(const sem::Variable* var) {
+bool GeneratorImpl::EmitIOVariable(const sem::GlobalVariable* var) {
     auto* decl = var->Declaration();
 
     if (auto* b = ast::GetAttribute<ast::BuiltinAttribute>(decl->attributes)) {
@@ -2018,7 +2018,7 @@
     }
 
     auto out = line();
-    EmitAttributes(out, decl->attributes);
+    EmitAttributes(out, var, decl->attributes);
     EmitInterpolationQualifiers(out, decl->attributes);
 
     auto name = builder_.Symbols().NameFor(decl->symbol);
@@ -2065,15 +2065,16 @@
 }
 
 bool GeneratorImpl::EmitAttributes(std::ostream& out,
+                                   const sem::GlobalVariable* var,
                                    utils::VectorRef<const ast::Attribute*> attributes) {
     if (attributes.IsEmpty()) {
         return true;
     }
     bool first = true;
     for (auto* attr : attributes) {
-        if (auto* location = attr->As<ast::LocationAttribute>()) {
+        if (attr->As<ast::LocationAttribute>()) {
             out << (first ? "layout(" : ", ");
-            out << "location = " << std::to_string(location->value);
+            out << "location = " << std::to_string(var->Location().value());
             first = false;
         }
     }
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index 502df8b..e70bdc2 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -324,7 +324,7 @@
     /// Handles emitting a global variable with the input or output storage class
     /// @param var the global variable
     /// @returns true on success
-    bool EmitIOVariable(const sem::Variable* var);
+    bool EmitIOVariable(const sem::GlobalVariable* var);
 
     /// Handles emitting interpolation qualifiers
     /// @param out the output of the expression stream
@@ -333,9 +333,12 @@
                                      utils::VectorRef<const ast::Attribute*> attrs);
     /// Handles emitting attributes
     /// @param out the output of the expression stream
+    /// @param var the global variable semantics
     /// @param attrs the attributes
     /// @returns true if the attributes were emitted
-    bool EmitAttributes(std::ostream& out, utils::VectorRef<const ast::Attribute*> attrs);
+    bool EmitAttributes(std::ostream& out,
+                        const sem::GlobalVariable* var,
+                        utils::VectorRef<const ast::Attribute*> attrs);
     /// Handles emitting the entry point function
     /// @param func the entry point
     /// @returns true if the entry point function was emitted
diff --git a/src/tint/writer/glsl/generator_impl_function_test.cc b/src/tint/writer/glsl/generator_impl_function_test.cc
index c388005..fd74e2d 100644
--- a/src/tint/writer/glsl/generator_impl_function_test.cc
+++ b/src/tint/writer/glsl/generator_impl_function_test.cc
@@ -128,7 +128,7 @@
     // }
     Func("frag_main",
          utils::Vector{
-             Param("foo", ty.f32(), utils::Vector{Location(0)}),
+             Param("foo", ty.f32(), utils::Vector{Location(0_a)}),
          },
          ty.f32(),
          utils::Vector{
@@ -138,7 +138,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(1),
+             Location(1_a),
          });
 
     GeneratorImpl& gen = SanitizeAndBuild();
@@ -218,8 +218,8 @@
         "Interface",
         utils::Vector{
             Member("pos", ty.vec4<f32>(), utils::Vector{Builtin(ast::BuiltinValue::kPosition)}),
-            Member("col1", ty.f32(), utils::Vector{Location(1)}),
-            Member("col2", ty.f32(), utils::Vector{Location(2)}),
+            Member("col1", ty.f32(), utils::Vector{Location(1_a)}),
+            Member("col2", ty.f32(), utils::Vector{Location(2_a)}),
         });
 
     Func("vert_main", utils::Empty, ty.Of(interface_struct),
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 3625db6..b05db3e 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -3947,23 +3947,24 @@
             std::string pre, post;
             if (auto* decl = mem->Declaration()) {
                 for (auto* attr : decl->attributes) {
-                    if (auto* location = attr->As<ast::LocationAttribute>()) {
+                    if (attr->Is<ast::LocationAttribute>()) {
                         auto& pipeline_stage_uses = str->PipelineStageUses();
                         if (pipeline_stage_uses.size() != 1) {
                             TINT_ICE(Writer, diagnostics_) << "invalid entry point IO struct uses";
                         }
 
+                        auto loc = mem->Location().value();
                         if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
-                            post += " : TEXCOORD" + std::to_string(location->value);
+                            post += " : TEXCOORD" + std::to_string(loc);
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kVertexOutput)) {
-                            post += " : TEXCOORD" + std::to_string(location->value);
+                            post += " : TEXCOORD" + std::to_string(loc);
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kFragmentInput)) {
-                            post += " : TEXCOORD" + std::to_string(location->value);
+                            post += " : TEXCOORD" + std::to_string(loc);
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kFragmentOutput)) {
-                            post += " : SV_Target" + std::to_string(location->value);
+                            post += " : SV_Target" + std::to_string(loc);
                         } else {
                             TINT_ICE(Writer, diagnostics_) << "invalid use of location attribute";
                         }
diff --git a/src/tint/writer/hlsl/generator_impl_function_test.cc b/src/tint/writer/hlsl/generator_impl_function_test.cc
index 14e8a70..bcd1891 100644
--- a/src/tint/writer/hlsl/generator_impl_function_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_function_test.cc
@@ -117,7 +117,7 @@
     // fn frag_main(@location(0) foo : f32) -> @location(1) f32 {
     //   return foo;
     // }
-    auto* foo_in = Param("foo", ty.f32(), utils::Vector{Location(0)});
+    auto* foo_in = Param("foo", ty.f32(), utils::Vector{Location(0_a)});
     Func("frag_main", utils::Vector{foo_in}, ty.f32(),
          utils::Vector{
              Return("foo"),
@@ -126,7 +126,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(1),
+             Location(1_a),
          });
 
     GeneratorImpl& gen = SanitizeAndBuild();
@@ -210,8 +210,8 @@
         "Interface",
         utils::Vector{
             Member("pos", ty.vec4<f32>(), utils::Vector{Builtin(ast::BuiltinValue::kPosition)}),
-            Member("col1", ty.f32(), utils::Vector{Location(1)}),
-            Member("col2", ty.f32(), utils::Vector{Location(2)}),
+            Member("col1", ty.f32(), utils::Vector{Location(1_a)}),
+            Member("col2", ty.f32(), utils::Vector{Location(2_a)}),
         });
 
     Func("vert_main", utils::Empty, ty.Of(interface_struct),
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 7617678..70bdb61 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -2785,24 +2785,25 @@
                         out << " [[" << name << "]]";
                         return true;
                     },
-                    [&](const ast::LocationAttribute* loc) {
+                    [&](const ast::LocationAttribute*) {
                         auto& pipeline_stage_uses = str->PipelineStageUses();
                         if (pipeline_stage_uses.size() != 1) {
                             TINT_ICE(Writer, diagnostics_) << "invalid entry point IO struct uses";
                             return false;
                         }
 
+                        uint32_t loc = mem->Location().value();
                         if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
-                            out << " [[attribute(" + std::to_string(loc->value) + ")]]";
+                            out << " [[attribute(" + std::to_string(loc) + ")]]";
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kVertexOutput)) {
-                            out << " [[user(locn" + std::to_string(loc->value) + ")]]";
+                            out << " [[user(locn" + std::to_string(loc) + ")]]";
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kFragmentInput)) {
-                            out << " [[user(locn" + std::to_string(loc->value) + ")]]";
+                            out << " [[user(locn" + std::to_string(loc) + ")]]";
                         } else if (pipeline_stage_uses.count(
                                        sem::PipelineStageUsage::kFragmentOutput)) {
-                            out << " [[color(" + std::to_string(loc->value) + ")]]";
+                            out << " [[color(" + std::to_string(loc) + ")]]";
                         } else {
                             TINT_ICE(Writer, diagnostics_) << "invalid use of location decoration";
                             return false;
diff --git a/src/tint/writer/msl/generator_impl_function_test.cc b/src/tint/writer/msl/generator_impl_function_test.cc
index fd612b3..addd255 100644
--- a/src/tint/writer/msl/generator_impl_function_test.cc
+++ b/src/tint/writer/msl/generator_impl_function_test.cc
@@ -91,7 +91,7 @@
     // fn frag_main(@location(0) foo : f32) -> @location(1) f32 {
     //   return foo;
     // }
-    auto* foo_in = Param("foo", ty.f32(), utils::Vector{Location(0)});
+    auto* foo_in = Param("foo", ty.f32(), utils::Vector{Location(0_a)});
     Func("frag_main", utils::Vector{foo_in}, ty.f32(),
          utils::Vector{
              Return("foo"),
@@ -100,7 +100,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(1),
+             Location(1_a),
          });
 
     GeneratorImpl& gen = SanitizeAndBuild();
@@ -188,8 +188,8 @@
     auto* interface_struct = Structure(
         "Interface",
         utils::Vector{
-            Member("col1", ty.f32(), utils::Vector{Location(1)}),
-            Member("col2", ty.f32(), utils::Vector{Location(2)}),
+            Member("col1", ty.f32(), utils::Vector{Location(1_a)}),
+            Member("col2", ty.f32(), utils::Vector{Location(2_a)}),
             Member("pos", ty.vec4<f32>(), utils::Vector{Builtin(ast::BuiltinValue::kPosition)}),
         });
 
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 04cbaac..3ee241d 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -884,9 +884,9 @@
                             U32Operand(ConvertBuiltin(builtin->builtin, sem->StorageClass()))});
                 return true;
             },
-            [&](const ast::LocationAttribute* location) {
+            [&](const ast::LocationAttribute*) {
                 push_annot(spv::Op::OpDecorate, {Operand(var_id), U32Operand(SpvDecorationLocation),
-                                                 Operand(location->value)});
+                                                 Operand(sem->Location().value())});
                 return true;
             },
             [&](const ast::InterpolateAttribute* interpolate) {
diff --git a/src/tint/writer/spirv/builder_binary_expression_test.cc b/src/tint/writer/spirv/builder_binary_expression_test.cc
index 9c4ba08..fd9a64a 100644
--- a/src/tint/writer/spirv/builder_binary_expression_test.cc
+++ b/src/tint/writer/spirv/builder_binary_expression_test.cc
@@ -981,36 +981,52 @@
 }
 
 TEST_F(BuilderTest, Binary_LogicalAnd) {
-    auto* lhs = create<ast::BinaryExpression>(ast::BinaryOp::kEqual, Expr(1_i), Expr(2_i));
-    auto* rhs = create<ast::BinaryExpression>(ast::BinaryOp::kEqual, Expr(3_i), Expr(4_i));
-    auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd, lhs, rhs);
+    auto* v0 = Var("a", Expr(1_i));
+    auto* v1 = Var("b", Expr(2_i));
+    auto* v2 = Var("c", Expr(3_i));
+    auto* v3 = Var("d", Expr(4_i));
+    auto* expr = LogicalAnd(Equal("a", "b"), Equal("c", "d"));
 
-    WrapInFunction(expr);
+    WrapInFunction(v0, v1, v2, v3, expr);
 
     spirv::Builder& b = Build();
 
     b.push_function(Function{});
     b.GenerateLabel(b.next_id());
+    ASSERT_TRUE(b.GenerateFunctionVariable(v0)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v1)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v2)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v3)) << b.error();
 
-    EXPECT_EQ(b.GenerateBinaryExpression(expr), 12u) << b.error();
+    EXPECT_EQ(b.GenerateBinaryExpression(expr), 22u) << b.error();
     EXPECT_EQ(DumpInstructions(b.types()),
               R"(%2 = OpTypeInt 32 1
 %3 = OpConstant %2 1
-%4 = OpConstant %2 2
-%6 = OpTypeBool
+%5 = OpTypePointer Function %2
+%6 = OpConstantNull %2
+%7 = OpConstant %2 2
 %9 = OpConstant %2 3
-%10 = OpConstant %2 4
+%11 = OpConstant %2 4
+%16 = OpTypeBool
 )");
     EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
               R"(%1 = OpLabel
-%5 = OpIEqual %6 %3 %4
-OpSelectionMerge %7 None
-OpBranchConditional %5 %8 %7
-%8 = OpLabel
-%11 = OpIEqual %6 %9 %10
-OpBranch %7
-%7 = OpLabel
-%12 = OpPhi %6 %5 %1 %11 %8
+OpStore %4 %3
+OpStore %8 %7
+OpStore %10 %9
+OpStore %12 %11
+%13 = OpLoad %2 %4
+%14 = OpLoad %2 %8
+%15 = OpIEqual %16 %13 %14
+OpSelectionMerge %17 None
+OpBranchConditional %15 %18 %17
+%18 = OpLabel
+%19 = OpLoad %2 %10
+%20 = OpLoad %2 %12
+%21 = OpIEqual %16 %19 %20
+OpBranch %17
+%17 = OpLabel
+%22 = OpPhi %16 %15 %1 %21 %18
 )");
 }
 
@@ -1131,38 +1147,52 @@
 }
 
 TEST_F(BuilderTest, Binary_LogicalOr) {
-    auto* lhs = create<ast::BinaryExpression>(ast::BinaryOp::kEqual, Expr(1_i), Expr(2_i));
+    auto* v0 = Var("a", Expr(1_i));
+    auto* v1 = Var("b", Expr(2_i));
+    auto* v2 = Var("c", Expr(3_i));
+    auto* v3 = Var("d", Expr(4_i));
+    auto* expr = LogicalOr(Equal("a", "b"), Equal("c", "d"));
 
-    auto* rhs = create<ast::BinaryExpression>(ast::BinaryOp::kEqual, Expr(3_i), Expr(4_i));
-
-    auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr, lhs, rhs);
-
-    WrapInFunction(expr);
+    WrapInFunction(v0, v1, v2, v3, expr);
 
     spirv::Builder& b = Build();
 
     b.push_function(Function{});
     b.GenerateLabel(b.next_id());
+    ASSERT_TRUE(b.GenerateFunctionVariable(v0)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v1)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v2)) << b.error();
+    ASSERT_TRUE(b.GenerateFunctionVariable(v3)) << b.error();
 
-    EXPECT_EQ(b.GenerateBinaryExpression(expr), 12u) << b.error();
+    EXPECT_EQ(b.GenerateBinaryExpression(expr), 22u) << b.error();
     EXPECT_EQ(DumpInstructions(b.types()),
               R"(%2 = OpTypeInt 32 1
 %3 = OpConstant %2 1
-%4 = OpConstant %2 2
-%6 = OpTypeBool
+%5 = OpTypePointer Function %2
+%6 = OpConstantNull %2
+%7 = OpConstant %2 2
 %9 = OpConstant %2 3
-%10 = OpConstant %2 4
+%11 = OpConstant %2 4
+%16 = OpTypeBool
 )");
     EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
               R"(%1 = OpLabel
-%5 = OpIEqual %6 %3 %4
-OpSelectionMerge %7 None
-OpBranchConditional %5 %7 %8
-%8 = OpLabel
-%11 = OpIEqual %6 %9 %10
-OpBranch %7
-%7 = OpLabel
-%12 = OpPhi %6 %5 %1 %11 %8
+OpStore %4 %3
+OpStore %8 %7
+OpStore %10 %9
+OpStore %12 %11
+%13 = OpLoad %2 %4
+%14 = OpLoad %2 %8
+%15 = OpIEqual %16 %13 %14
+OpSelectionMerge %17 None
+OpBranchConditional %15 %17 %18
+%18 = OpLabel
+%19 = OpLoad %2 %10
+%20 = OpLoad %2 %12
+%21 = OpIEqual %16 %19 %20
+OpBranch %17
+%17 = OpLabel
+%22 = OpPhi %16 %15 %1 %21 %18
 )");
 }
 
diff --git a/src/tint/writer/spirv/builder_entry_point_test.cc b/src/tint/writer/spirv/builder_entry_point_test.cc
index 424e89b..a4128c1 100644
--- a/src/tint/writer/spirv/builder_entry_point_test.cc
+++ b/src/tint/writer/spirv/builder_entry_point_test.cc
@@ -48,7 +48,7 @@
                         });
     auto* loc1 = Param("loc1", ty.f32(),
                        utils::Vector{
-                           Location(1u),
+                           Location(1_u),
                        });
     auto* mul = Mul(Expr(MemberAccessor("coord", "x")), Expr("loc1"));
     auto* col = Var("col", ty.f32(), mul);
@@ -120,7 +120,7 @@
     // }
     auto* loc_in = Param("loc_in", ty.u32(),
                          utils::Vector{
-                             Location(0),
+                             Location(0_a),
                              Flat(),
                          });
     auto* cond =
@@ -134,7 +134,7 @@
              Stage(ast::PipelineStage::kFragment),
          },
          utils::Vector{
-             Location(0),
+             Location(0_a),
          });
 
     spirv::Builder& b = SanitizeAndBuild();
@@ -211,7 +211,7 @@
     auto* interface = Structure(
         "Interface",
         utils::Vector{
-            Member("value", ty.f32(), utils::Vector{Location(1u)}),
+            Member("value", ty.f32(), utils::Vector{Location(1_u)}),
             Member("pos", ty.vec4<f32>(), utils::Vector{Builtin(ast::BuiltinValue::kPosition)}),
         });
 
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index bd75fb8..b2fba29 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -741,7 +741,7 @@
             },
             [&](const ast::BindingAttribute* binding) {
                 out << "binding(";
-                if (!EmitExpression(out, binding->value)) {
+                if (!EmitExpression(out, binding->expr)) {
                     return false;
                 }
                 out << ")";
@@ -749,14 +749,18 @@
             },
             [&](const ast::GroupAttribute* group) {
                 out << "group(";
-                if (!EmitExpression(out, group->value)) {
+                if (!EmitExpression(out, group->expr)) {
                     return false;
                 }
                 out << ")";
                 return true;
             },
             [&](const ast::LocationAttribute* location) {
-                out << "location(" << location->value << ")";
+                out << "location(";
+                if (!EmitExpression(out, location->expr)) {
+                    return false;
+                }
+                out << ")";
                 return true;
             },
             [&](const ast::BuiltinAttribute* builtin) {
@@ -777,7 +781,7 @@
             },
             [&](const ast::IdAttribute* override_deco) {
                 out << "id(";
-                if (!EmitExpression(out, override_deco->value)) {
+                if (!EmitExpression(out, override_deco->expr)) {
                     return false;
                 }
                 out << ")";
@@ -789,7 +793,7 @@
             },
             [&](const ast::StructMemberAlignAttribute* align) {
                 out << "align(";
-                if (!EmitExpression(out, align->align)) {
+                if (!EmitExpression(out, align->expr)) {
                     return false;
                 }
                 out << ")";
diff --git a/src/tint/writer/wgsl/generator_impl_function_test.cc b/src/tint/writer/wgsl/generator_impl_function_test.cc
index 3b80e69..af09a13 100644
--- a/src/tint/writer/wgsl/generator_impl_function_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_function_test.cc
@@ -116,7 +116,7 @@
                         });
     auto* loc1 = Param("loc1", ty.f32(),
                        utils::Vector{
-                           Location(1u),
+                           Location(1_a),
                        });
     auto* func = Func("frag_main", utils::Vector{coord, loc1}, ty.void_(), utils::Empty,
                       utils::Vector{
@@ -143,7 +143,7 @@
                           Stage(ast::PipelineStage::kFragment),
                       },
                       utils::Vector{
-                          Location(1u),
+                          Location(1_a),
                       });
 
     GeneratorImpl& gen = Build();
diff --git a/src/tint/writer/wgsl/generator_impl_type_test.cc b/src/tint/writer/wgsl/generator_impl_type_test.cc
index 5390057..ef90579 100644
--- a/src/tint/writer/wgsl/generator_impl_type_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_type_test.cc
@@ -274,7 +274,7 @@
     auto* s = Structure(
         "S", utils::Vector{
                  Member("a", ty.u32(), utils::Vector{Builtin(ast::BuiltinValue::kVertexIndex)}),
-                 Member("b", ty.f32(), utils::Vector{Location(2u)}),
+                 Member("b", ty.f32(), utils::Vector{Location(2_a)}),
              });
 
     GeneratorImpl& gen = Build();