From 0f2d95dea31e4e502071c6487f95c247392c93c2 Mon Sep 17 00:00:00 2001
From: Ben Clayton <bclayton@google.com>
Date: Thu, 22 Jul 2021 13:24:59 +0000
Subject: [PATCH] sem: Split sem::Variable into global, local and parameter

Each of these may contain information specific to their kind.

Change-Id: Ic8ac808088132b7bc2e43da6ce46a06571e0fed5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59200
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
---
 src/ast/identifier_expression.h               |   1 -
 src/inspector/inspector.cc                    |  15 +-
 src/inspector/inspector_test.cc               |  15 +-
 src/intrinsic_table.cc                        |  80 +++++++++-
 src/intrinsic_table_test.cc                   | 143 +++++++++++-------
 src/program_builder.h                         |   1 +
 src/reader/spirv/function.cc                  |   5 +-
 src/resolver/intrinsic_test.cc                |   4 +-
 .../pipeline_overridable_constant_test.cc     |   8 +-
 src/resolver/resolver.cc                      |  29 +++-
 src/resolver/resolver_test.cc                 |   6 +-
 src/sem/call_target.cc                        |   8 +-
 src/sem/call_target.h                         |  41 +----
 src/sem/function.cc                           |  18 +--
 src/sem/function.h                            |   6 +-
 src/sem/intrinsic.cc                          |  29 +---
 src/sem/intrinsic.h                           |  17 +--
 src/sem/intrinsic_type.cc                     |   7 +
 src/sem/intrinsic_type.cc.tmpl                |   7 +
 src/sem/intrinsic_type.h                      |   5 +
 src/sem/intrinsic_type.h.tmpl                 |   5 +
 src/sem/sampler_type_test.cc                  |   1 +
 src/sem/variable.cc                           |  50 ++++--
 src/sem/variable.h                            | 119 ++++++++++++---
 src/transform/array_length_from_uniform.cc    |   8 +-
 src/transform/decompose_memory_access.cc      |   4 +-
 src/transform/msl.cc                          |   6 +-
 src/transform/robustness.cc                   |   2 +-
 src/transform/transform.cc                    |   1 +
 src/utils/hash.h                              |  15 +-
 src/writer/hlsl/generator_impl.cc             |  24 +--
 src/writer/hlsl/generator_impl_type_test.cc   |   1 +
 src/writer/msl/generator_impl.cc              |   6 +-
 src/writer/msl/generator_impl_type_test.cc    |   1 +
 src/writer/spirv/builder.cc                   |  23 +--
 src/writer/spirv/builder.h                    |   1 +
 36 files changed, 435 insertions(+), 277 deletions(-)

diff --git a/src/ast/identifier_expression.h b/src/ast/identifier_expression.h
index da32342c44..3deff16857 100644
--- a/src/ast/identifier_expression.h
+++ b/src/ast/identifier_expression.h
@@ -16,7 +16,6 @@
 #define SRC_AST_IDENTIFIER_EXPRESSION_H_
 
 #include "src/ast/expression.h"
-#include "src/sem/intrinsic.h"
 
 namespace tint {
 namespace ast {
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index 1003e0faff..0e64db75c4 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -174,7 +174,8 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
 
       auto name = program_->Symbols().NameFor(decl->symbol());
 
-      if (var->IsPipelineConstant()) {
+      auto* global = var->As<sem::GlobalVariable>();
+      if (global && global->IsPipelineConstant()) {
         OverridableConstant overridable_constant;
         overridable_constant.name = name;
         entry_point.overridable_constants.push_back(overridable_constant);
@@ -203,8 +204,8 @@ std::string Inspector::GetRemappedNameForEntryPoint(
 std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
   std::map<uint32_t, Scalar> result;
   for (auto* var : program_->AST().GlobalVariables()) {
-    auto* sem_var = program_->Sem().Get(var);
-    if (!sem_var->IsPipelineConstant()) {
+    auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
+    if (!global || !global->IsPipelineConstant()) {
       continue;
     }
 
@@ -212,7 +213,7 @@ std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
     // WGSL, so the resolver should catch it. Thus here the inspector just
     // assumes all definitions of the constant id are the same, so only needs
     // to find the first reference to constant id.
-    uint32_t constant_id = sem_var->ConstantId();
+    uint32_t constant_id = global->ConstantId();
     if (result.find(constant_id) != result.end()) {
       continue;
     }
@@ -274,10 +275,10 @@ std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
 std::map<std::string, uint32_t> Inspector::GetConstantNameToIdMap() {
   std::map<std::string, uint32_t> result;
   for (auto* var : program_->AST().GlobalVariables()) {
-    auto* sem_var = program_->Sem().Get(var);
-    if (sem_var->IsPipelineConstant()) {
+    auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
+    if (global && global->IsPipelineConstant()) {
       auto name = program_->Symbols().NameFor(var->symbol());
-      result[name] = sem_var->ConstantId();
+      result[name] = global->ConstantId();
     }
   }
   return result;
diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc
index 7d7740a437..40fc59147c 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -928,16 +928,19 @@ TEST_F(InspectorGetConstantNameToIdMapTest, WithAndWithoutIds) {
   EXPECT_EQ(result["v300"], 300u);
 
   ASSERT_TRUE(result.count("a"));
-  ASSERT_TRUE(program_->Sem().Get(a));
-  EXPECT_EQ(result["a"], program_->Sem().Get(a)->ConstantId());
+  ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(a));
+  EXPECT_EQ(result["a"],
+            program_->Sem().Get<sem::GlobalVariable>(a)->ConstantId());
 
   ASSERT_TRUE(result.count("b"));
-  ASSERT_TRUE(program_->Sem().Get(b));
-  EXPECT_EQ(result["b"], program_->Sem().Get(b)->ConstantId());
+  ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(b));
+  EXPECT_EQ(result["b"],
+            program_->Sem().Get<sem::GlobalVariable>(b)->ConstantId());
 
   ASSERT_TRUE(result.count("c"));
-  ASSERT_TRUE(program_->Sem().Get(c));
-  EXPECT_EQ(result["c"], program_->Sem().Get(c)->ConstantId());
+  ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(c));
+  EXPECT_EQ(result["c"],
+            program_->Sem().Get<sem::GlobalVariable>(c)->ConstantId());
 }
 
 TEST_F(InspectorGetStorageSizeTest, Empty) {
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc
index 233aff4a8e..4f1fa28276 100644
--- a/src/intrinsic_table.cc
+++ b/src/intrinsic_table.cc
@@ -701,6 +701,55 @@ struct IntrinsicInfo {
 
 #include "intrinsic_table.inl"
 
+/// IntrinsicPrototype describes a fully matched intrinsic function, which is
+/// used as a lookup for building unique sem::Intrinsic instances.
+struct IntrinsicPrototype {
+  /// Parameter describes a single parameter
+  struct Parameter {
+    /// Parameter type
+    sem::Type* const type;
+    /// Parameter usage
+    ParameterUsage const usage = ParameterUsage::kNone;
+  };
+
+  /// Hasher provides a hash function for the IntrinsicPrototype
+  struct Hasher {
+    /// @param i the IntrinsicPrototype to create a hash for
+    /// @return the hash value
+    inline std::size_t operator()(const IntrinsicPrototype& i) const {
+      size_t hash = utils::Hash(i.parameters.size());
+      for (auto& p : i.parameters) {
+        utils::HashCombine(&hash, p.type, p.usage);
+      }
+      return utils::Hash(hash, i.type, i.return_type, i.supported_stages,
+                         i.is_deprecated);
+    }
+  };
+
+  sem::IntrinsicType type = sem::IntrinsicType::kNone;
+  std::vector<Parameter> parameters;
+  sem::Type const* return_type = nullptr;
+  PipelineStageSet supported_stages;
+  bool is_deprecated = false;
+};
+
+/// Equality operator for IntrinsicPrototype
+bool operator==(const IntrinsicPrototype& a, const IntrinsicPrototype& b) {
+  if (a.type != b.type || a.supported_stages != b.supported_stages ||
+      a.return_type != b.return_type || a.is_deprecated != b.is_deprecated ||
+      a.parameters.size() != b.parameters.size()) {
+    return false;
+  }
+  for (size_t i = 0; i < a.parameters.size(); i++) {
+    auto& pa = a.parameters[i];
+    auto& pb = b.parameters[i];
+    if (pa.type != pb.type || pa.usage != pb.usage) {
+      return false;
+    }
+  }
+  return true;
+}
+
 /// Impl is the private implementation of the IntrinsicTable interface.
 class Impl : public IntrinsicTable {
  public:
@@ -726,7 +775,10 @@ class Impl : public IntrinsicTable {
 
   ProgramBuilder& builder;
   Matchers matchers;
-  std::unordered_map<sem::Intrinsic, sem::Intrinsic*> intrinsics;
+  std::unordered_map<IntrinsicPrototype,
+                     sem::Intrinsic*,
+                     IntrinsicPrototype::Hasher>
+      intrinsics;
 };
 
 /// @return a string representing a call to an intrinsic with the given argument
@@ -833,7 +885,7 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
 
   ClosedState closed(builder);
 
-  sem::ParameterList parameters;
+  std::vector<IntrinsicPrototype::Parameter> parameters;
 
   auto num_params = std::min(num_parameters, num_arguments);
   for (uint32_t p = 0; p < num_params; p++) {
@@ -841,8 +893,8 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
     auto* indices = parameter.matcher_indices;
     auto* type = Match(closed, overload, indices).Type(args[p]->UnwrapRef());
     if (type) {
-      parameters.emplace_back(
-          sem::Parameter{const_cast<sem::Type*>(type), parameter.usage});
+      parameters.emplace_back(IntrinsicPrototype::Parameter{
+          const_cast<sem::Type*>(type), parameter.usage});
       match_score += kScorePerMatchedParam;
     } else {
       overload_matched = false;
@@ -899,13 +951,25 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
     return_type = builder.create<sem::Void>();
   }
 
-  sem::Intrinsic intrinsic(intrinsic_type, const_cast<sem::Type*>(return_type),
-                           std::move(parameters), overload.supported_stages,
-                           overload.is_deprecated);
+  IntrinsicPrototype intrinsic;
+  intrinsic.type = intrinsic_type;
+  intrinsic.return_type = return_type;
+  intrinsic.parameters = std::move(parameters);
+  intrinsic.supported_stages = overload.supported_stages;
+  intrinsic.is_deprecated = overload.is_deprecated;
 
   // De-duplicate intrinsics that are identical.
   return utils::GetOrCreate(intrinsics, intrinsic, [&] {
-    return builder.create<sem::Intrinsic>(intrinsic);
+    sem::ParameterList params;
+    params.reserve(intrinsic.parameters.size());
+    for (auto& p : intrinsic.parameters) {
+      params.emplace_back(builder.create<sem::Parameter>(
+          nullptr, p.type, ast::StorageClass::kNone, ast::Access::kUndefined,
+          p.usage));
+    }
+    return builder.create<sem::Intrinsic>(
+        intrinsic.type, const_cast<sem::Type*>(intrinsic.return_type),
+        std::move(params), intrinsic.supported_stages, intrinsic.is_deprecated);
   });
 }
 
diff --git a/src/intrinsic_table_test.cc b/src/intrinsic_table_test.cc
index 2635312a4e..374c56b895 100644
--- a/src/intrinsic_table_test.cc
+++ b/src/intrinsic_table_test.cc
@@ -45,7 +45,8 @@ TEST_F(IntrinsicTableTest, MatchF32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kCos);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{f32}));
+  ASSERT_THAT(result->Parameters().size(), 1);
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchF32) {
@@ -65,7 +66,8 @@ TEST_F(IntrinsicTableTest, MatchU32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kUnpack2x16float);
   EXPECT_THAT(result->ReturnType(), vec2_f32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{u32}));
+  ASSERT_EQ(result->Parameters().size(), 1u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), u32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchU32) {
@@ -87,10 +89,13 @@ TEST_F(IntrinsicTableTest, MatchI32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{i32, ParameterUsage::kCoords},
-                          Parameter{i32, ParameterUsage::kLevel}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
+  EXPECT_EQ(result->Parameters()[2]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
 }
 
 TEST_F(IntrinsicTableTest, MismatchI32) {
@@ -109,7 +114,8 @@ TEST_F(IntrinsicTableTest, MatchIU32AsI32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kCountOneBits);
   EXPECT_THAT(result->ReturnType(), i32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{i32}));
+  ASSERT_EQ(result->Parameters().size(), 1u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), i32);
 }
 
 TEST_F(IntrinsicTableTest, MatchIU32AsU32) {
@@ -119,7 +125,8 @@ TEST_F(IntrinsicTableTest, MatchIU32AsU32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kCountOneBits);
   EXPECT_THAT(result->ReturnType(), u32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{u32}));
+  ASSERT_EQ(result->Parameters().size(), 1u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), u32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchIU32) {
@@ -137,8 +144,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsI32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
   EXPECT_THAT(result->ReturnType(), i32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{i32}, Parameter{i32}, Parameter{i32}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), i32);
 }
 
 TEST_F(IntrinsicTableTest, MatchFIU32AsU32) {
@@ -149,8 +158,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsU32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
   EXPECT_THAT(result->ReturnType(), u32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{u32}, Parameter{u32}, Parameter{u32}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), u32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), u32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), u32);
 }
 
 TEST_F(IntrinsicTableTest, MatchFIU32AsF32) {
@@ -161,8 +172,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsF32) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{f32}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), f32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchFIU32) {
@@ -182,8 +195,10 @@ TEST_F(IntrinsicTableTest, MatchBool) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kSelect);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{bool_}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), bool_);
 }
 
 TEST_F(IntrinsicTableTest, MismatchBool) {
@@ -203,8 +218,9 @@ TEST_F(IntrinsicTableTest, MatchPointer) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kModf);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{f32}, Parameter{ptr}));
+  ASSERT_EQ(result->Parameters().size(), 2u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), ptr);
 }
 
 TEST_F(IntrinsicTableTest, MismatchPointer) {
@@ -225,7 +241,7 @@ TEST_F(IntrinsicTableTest, MatchArray) {
   EXPECT_THAT(result->Type(), IntrinsicType::kArrayLength);
   EXPECT_TRUE(result->ReturnType()->Is<sem::U32>());
   ASSERT_EQ(result->Parameters().size(), 1u);
-  auto* param_type = result->Parameters()[0].type;
+  auto* param_type = result->Parameters()[0]->Type();
   ASSERT_TRUE(param_type->Is<sem::Pointer>());
   EXPECT_TRUE(param_type->As<sem::Pointer>()->StoreType()->Is<sem::Array>());
 }
@@ -249,10 +265,13 @@ TEST_F(IntrinsicTableTest, MatchSampler) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureSample);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{sampler, ParameterUsage::kSampler},
-                          Parameter{vec2_f32, ParameterUsage::kCoords}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), sampler);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kSampler);
+  EXPECT_EQ(result->Parameters()[2]->Type(), vec2_f32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kCoords);
 }
 
 TEST_F(IntrinsicTableTest, MismatchSampler) {
@@ -277,10 +296,13 @@ TEST_F(IntrinsicTableTest, MatchSampledTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords},
-                          Parameter{i32, ParameterUsage::kLevel}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
+  EXPECT_EQ(result->Parameters()[2]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
 }
 
 TEST_F(IntrinsicTableTest, MatchMultisampledTexture) {
@@ -295,10 +317,13 @@ TEST_F(IntrinsicTableTest, MatchMultisampledTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords},
-                          Parameter{i32, ParameterUsage::kSampleIndex}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
+  EXPECT_EQ(result->Parameters()[2]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kSampleIndex);
 }
 
 TEST_F(IntrinsicTableTest, MatchDepthTexture) {
@@ -312,10 +337,13 @@ TEST_F(IntrinsicTableTest, MatchDepthTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords},
-                          Parameter{i32, ParameterUsage::kLevel}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
+  EXPECT_EQ(result->Parameters()[2]->Type(), i32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
 }
 
 TEST_F(IntrinsicTableTest, MatchExternalTexture) {
@@ -330,9 +358,11 @@ TEST_F(IntrinsicTableTest, MatchExternalTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords}));
+  ASSERT_EQ(result->Parameters().size(), 2u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
 }
 
 TEST_F(IntrinsicTableTest, MatchROStorageTexture) {
@@ -352,9 +382,11 @@ TEST_F(IntrinsicTableTest, MatchROStorageTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
   EXPECT_THAT(result->ReturnType(), vec4_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords}));
+  ASSERT_EQ(result->Parameters().size(), 2u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
 }
 
 TEST_F(IntrinsicTableTest, MatchWOStorageTexture) {
@@ -374,10 +406,13 @@ TEST_F(IntrinsicTableTest, MatchWOStorageTexture) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kTextureStore);
   EXPECT_TRUE(result->ReturnType()->Is<sem::Void>());
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{tex, ParameterUsage::kTexture},
-                          Parameter{vec2_i32, ParameterUsage::kCoords},
-                          Parameter{vec4_f32, ParameterUsage::kValue}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), tex);
+  EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
+  EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
+  EXPECT_EQ(result->Parameters()[2]->Type(), vec4_f32);
+  EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kValue);
 }
 
 TEST_F(IntrinsicTableTest, MismatchTexture) {
@@ -401,7 +436,8 @@ TEST_F(IntrinsicTableTest, ImplicitLoadOnReference) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kCos);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{f32}));
+  ASSERT_EQ(result->Parameters().size(), 1u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
 }
 
 TEST_F(IntrinsicTableTest, MatchOpenType) {
@@ -412,8 +448,9 @@ TEST_F(IntrinsicTableTest, MatchOpenType) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{f32}));
+  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), f32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchOpenType) {
@@ -434,9 +471,10 @@ TEST_F(IntrinsicTableTest, MatchOpenSizeVector) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
   EXPECT_THAT(result->ReturnType(), vec2_f32);
-  EXPECT_THAT(result->Parameters(),
-              ElementsAre(Parameter{vec2_f32}, Parameter{vec2_f32},
-                          Parameter{vec2_f32}));
+  ASSERT_EQ(result->Parameters().size(), 3u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), vec2_f32);
+  EXPECT_EQ(result->Parameters()[1]->Type(), vec2_f32);
+  EXPECT_EQ(result->Parameters()[2]->Type(), vec2_f32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchOpenSizeVector) {
@@ -459,7 +497,8 @@ TEST_F(IntrinsicTableTest, MatchOpenSizeMatrix) {
   ASSERT_EQ(Diagnostics().str(), "");
   EXPECT_THAT(result->Type(), IntrinsicType::kDeterminant);
   EXPECT_THAT(result->ReturnType(), f32);
-  EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{mat3_f32}));
+  ASSERT_EQ(result->Parameters().size(), 1u);
+  EXPECT_EQ(result->Parameters()[0]->Type(), mat3_f32);
 }
 
 TEST_F(IntrinsicTableTest, MismatchOpenSizeMatrix) {
diff --git a/src/program_builder.h b/src/program_builder.h
index d4a25cfe52..a7d8f90aa9 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -50,6 +50,7 @@
 #include "src/ast/pointer.h"
 #include "src/ast/return_statement.h"
 #include "src/ast/sampled_texture.h"
+#include "src/ast/sampler.h"
 #include "src/ast/scalar_constructor_expression.h"
 #include "src/ast/sint_literal.h"
 #include "src/ast/stage_decoration.h"
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index d673c60bb6..819b9a30a3 100644
--- a/src/reader/spirv/function.cc
+++ b/src/reader/spirv/function.cc
@@ -34,6 +34,7 @@
 #include "src/ast/unary_op_expression.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/sem/depth_texture_type.h"
+#include "src/sem/intrinsic_type.h"
 #include "src/sem/sampled_texture_type.h"
 
 // Terms:
@@ -4959,9 +4960,7 @@ bool FunctionEmitter::EmitControlBarrier(
 TypedExpression FunctionEmitter::MakeIntrinsicCall(
     const spvtools::opt::Instruction& inst) {
   const auto intrinsic = GetIntrinsic(inst.opcode());
-  std::ostringstream ss;
-  ss << intrinsic;
-  auto name = ss.str();
+  auto* name = sem::str(intrinsic);
   auto* ident = create<ast::IdentifierExpression>(
       Source{}, builder_.Symbols().Register(name));
 
diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc
index 27b15df5b7..d94ed2d830 100644
--- a/src/resolver/intrinsic_test.cc
+++ b/src/resolver/intrinsic_test.cc
@@ -1706,11 +1706,11 @@ std::string to_str(const std::string& function,
   std::stringstream out;
   out << function << "(";
   bool first = true;
-  for (auto& param : params) {
+  for (auto* param : params) {
     if (!first) {
       out << ", ";
     }
-    out << sem::str(param.usage);
+    out << sem::str(param->Usage());
     first = false;
   }
   out << ")";
diff --git a/src/resolver/pipeline_overridable_constant_test.cc b/src/resolver/pipeline_overridable_constant_test.cc
index 9fdfd8fc30..5d0571f71e 100644
--- a/src/resolver/pipeline_overridable_constant_test.cc
+++ b/src/resolver/pipeline_overridable_constant_test.cc
@@ -30,7 +30,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, NonOverridable) {
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
-  auto* sem_a = Sem().Get(a);
+  auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
   ASSERT_NE(sem_a, nullptr);
   EXPECT_EQ(sem_a->Declaration(), a);
   EXPECT_FALSE(sem_a->IsPipelineConstant());
@@ -41,7 +41,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithId) {
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
-  auto* sem_a = Sem().Get(a);
+  auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
   ASSERT_NE(sem_a, nullptr);
   EXPECT_EQ(sem_a->Declaration(), a);
   EXPECT_TRUE(sem_a->IsPipelineConstant());
@@ -53,7 +53,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithoutId) {
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
-  auto* sem_a = Sem().Get(a);
+  auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
   ASSERT_NE(sem_a, nullptr);
   EXPECT_EQ(sem_a->Declaration(), a);
   EXPECT_TRUE(sem_a->IsPipelineConstant());
@@ -79,7 +79,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithAndWithoutIds) {
 
   std::vector<uint16_t> constant_ids;
   for (auto* var : variables) {
-    auto* sem = Sem().Get(var);
+    auto* sem = Sem().Get<sem::GlobalVariable>(var);
     ASSERT_NE(sem, nullptr);
     constant_ids.push_back(static_cast<uint16_t>(sem->ConstantId()));
   }
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index b41b1bcf69..842b34882f 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -3693,11 +3693,24 @@ void Resolver::CreateSemanticNodes() const {
         next_constant_id = constant_id + 1;
       }
 
-      sem_var = builder_->create<sem::Variable>(var, info->type, constant_id);
-    } else {
       sem_var =
-          builder_->create<sem::Variable>(var, info->type, info->storage_class,
-                                          info->access, info->binding_point);
+          builder_->create<sem::GlobalVariable>(var, info->type, constant_id);
+    } else {
+      switch (info->kind) {
+        case VariableKind::kGlobal:
+          sem_var = builder_->create<sem::GlobalVariable>(
+              var, info->type, info->storage_class, info->access,
+              info->binding_point);
+          break;
+        case VariableKind::kLocal:
+          sem_var = builder_->create<sem::LocalVariable>(
+              var, info->type, info->storage_class, info->access);
+          break;
+        case VariableKind::kParameter:
+          sem_var = builder_->create<sem::Parameter>(
+              var, info->type, info->storage_class, info->access);
+          break;
+      }
     }
 
     std::vector<const sem::VariableUser*> users;
@@ -3739,9 +3752,15 @@ void Resolver::CreateSemanticNodes() const {
     auto* func = it.first;
     auto* info = it.second;
 
+    sem::ParameterList parameters;
+    parameters.reserve(info->parameters.size());
+    for (auto* p : info->parameters) {
+      parameters.emplace_back(sem.Get<sem::Parameter>(p->declaration));
+    }
+
     auto* sem_func = builder_->create<sem::Function>(
         info->declaration, const_cast<sem::Type*>(info->return_type),
-        remap_vars(info->parameters), remap_vars(info->referenced_module_vars),
+        parameters, remap_vars(info->referenced_module_vars),
         remap_vars(info->local_referenced_module_vars), info->return_statements,
         info->callsites, ancestor_entry_points[func->symbol()],
         info->workgroup_size);
diff --git a/src/resolver/resolver_test.cc b/src/resolver/resolver_test.cc
index d231774ac1..f176dc190d 100644
--- a/src/resolver/resolver_test.cc
+++ b/src/resolver/resolver_test.cc
@@ -2018,8 +2018,10 @@ TEST_F(ResolverTest, BindingPoint_SetForResources) {
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
-  EXPECT_EQ(Sem().Get(s1)->BindingPoint(), (sem::BindingPoint{1u, 2u}));
-  EXPECT_EQ(Sem().Get(s2)->BindingPoint(), (sem::BindingPoint{3u, 4u}));
+  EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s1)->BindingPoint(),
+            (sem::BindingPoint{1u, 2u}));
+  EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s2)->BindingPoint(),
+            (sem::BindingPoint{3u, 4u}));
 }
 
 TEST_F(ResolverTest, Function_EntryPoints_StageDecoration) {
diff --git a/src/sem/call_target.cc b/src/sem/call_target.cc
index 5306bf8c52..b8bf12ca0d 100644
--- a/src/sem/call_target.cc
+++ b/src/sem/call_target.cc
@@ -32,18 +32,12 @@ CallTarget::~CallTarget() = default;
 
 int IndexOf(const ParameterList& parameters, ParameterUsage usage) {
   for (size_t i = 0; i < parameters.size(); i++) {
-    if (parameters[i].usage == usage) {
+    if (parameters[i]->Usage() == usage) {
       return static_cast<int>(i);
     }
   }
   return -1;
 }
 
-std::ostream& operator<<(std::ostream& out, Parameter parameter) {
-  out << "[type: " << parameter.type->FriendlyName(SymbolTable{ProgramID{}})
-      << ", usage: " << str(parameter.usage) << "]";
-  return out;
-}
-
 }  // namespace sem
 }  // namespace tint
diff --git a/src/sem/call_target.h b/src/sem/call_target.h
index e994c80f49..33d5029234 100644
--- a/src/sem/call_target.h
+++ b/src/sem/call_target.h
@@ -18,39 +18,15 @@
 #include <vector>
 
 #include "src/sem/node.h"
-#include "src/sem/parameter_usage.h"
 #include "src/sem/sampler_type.h"
+#include "src/sem/variable.h"
 #include "src/utils/hash.h"
 
 namespace tint {
-
 namespace sem {
 // Forward declarations
 class Type;
 
-/// Parameter describes a single parameter of a call target
-struct Parameter {
-  /// Parameter type
-  sem::Type* const type;
-  /// Parameter usage
-  ParameterUsage const usage = ParameterUsage::kNone;
-};
-
-std::ostream& operator<<(std::ostream& out, Parameter parameter);
-
-/// Equality operator for Parameters
-static inline bool operator==(const Parameter& a, const Parameter& b) {
-  return a.type == b.type && a.usage == b.usage;
-}
-
-/// Inequality operator for Parameters
-static inline bool operator!=(const Parameter& a, const Parameter& b) {
-  return !(a == b);
-}
-
-/// ParameterList is a list of Parameter
-using ParameterList = std::vector<Parameter>;
-
 /// @param parameters the list of parameters
 /// @param usage the parameter usage to find
 /// @returns the index of the parameter with the given usage, or -1 if no
@@ -85,19 +61,4 @@ class CallTarget : public Castable<CallTarget, Node> {
 }  // namespace sem
 }  // namespace tint
 
-namespace std {
-
-/// Custom std::hash specialization for tint::sem::Parameter
-template <>
-class hash<tint::sem::Parameter> {
- public:
-  /// @param p the tint::sem::Parameter to create a hash for
-  /// @return the hash value
-  inline std::size_t operator()(const tint::sem::Parameter& p) const {
-    return tint::utils::Hash(p.type, p.usage);
-  }
-};
-
-}  // namespace std
-
 #endif  // SRC_SEM_CALL_TARGET_H_
diff --git a/src/sem/function.cc b/src/sem/function.cc
index fcf192aee5..21fabf9490 100644
--- a/src/sem/function.cc
+++ b/src/sem/function.cc
@@ -27,31 +27,17 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Function);
 namespace tint {
 namespace sem {
 
-namespace {
-
-ParameterList GetParameters(const std::vector<const Variable*>& params) {
-  ParameterList parameters;
-  parameters.reserve(params.size());
-  for (auto* param : params) {
-    parameters.emplace_back(Parameter{param->Type(), ParameterUsage::kNone});
-  }
-  return parameters;
-}
-
-}  // namespace
-
 Function::Function(ast::Function* declaration,
                    Type* return_type,
-                   std::vector<const Variable*> parameters,
+                   ParameterList parameters,
                    std::vector<const Variable*> referenced_module_vars,
                    std::vector<const Variable*> local_referenced_module_vars,
                    std::vector<const ast::ReturnStatement*> return_statements,
                    std::vector<const ast::CallExpression*> callsites,
                    std::vector<Symbol> ancestor_entry_points,
                    std::array<WorkgroupDimension, 3> workgroup_size)
-    : Base(return_type, GetParameters(parameters)),
+    : Base(return_type, std::move(parameters)),
       declaration_(declaration),
-      parameters_(std::move(parameters)),
       referenced_module_vars_(std::move(referenced_module_vars)),
       local_referenced_module_vars_(std::move(local_referenced_module_vars)),
       return_statements_(std::move(return_statements)),
diff --git a/src/sem/function.h b/src/sem/function.h
index d48047b359..6ab1d39543 100644
--- a/src/sem/function.h
+++ b/src/sem/function.h
@@ -68,7 +68,7 @@ class Function : public Castable<Function, CallTarget> {
   /// @param workgroup_size the workgroup size
   Function(ast::Function* declaration,
            Type* return_type,
-           std::vector<const Variable*> parameters,
+           ParameterList parameters,
            std::vector<const Variable*> referenced_module_vars,
            std::vector<const Variable*> local_referenced_module_vars,
            std::vector<const ast::ReturnStatement*> return_statements,
@@ -82,9 +82,6 @@ class Function : public Castable<Function, CallTarget> {
   /// @returns the ast::Function declaration
   ast::Function* Declaration() const { return declaration_; }
 
-  /// @return the parameters to the function
-  const std::vector<const Variable*> Parameters() const { return parameters_; }
-
   /// Note: If this function calls other functions, the return will also include
   /// all of the referenced variables from the callees.
   /// @returns the referenced module variables
@@ -178,7 +175,6 @@ class Function : public Castable<Function, CallTarget> {
       bool multisampled) const;
 
   ast::Function* const declaration_;
-  std::vector<const Variable*> const parameters_;
   std::vector<const Variable*> const referenced_module_vars_;
   std::vector<const Variable*> const local_referenced_module_vars_;
   std::vector<const ast::ReturnStatement*> const return_statements_;
diff --git a/src/sem/intrinsic.cc b/src/sem/intrinsic.cc
index 4be75e8121..b7d57a932f 100644
--- a/src/sem/intrinsic.cc
+++ b/src/sem/intrinsic.cc
@@ -19,11 +19,6 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Intrinsic);
 namespace tint {
 namespace sem {
 
-std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
-  out << str(i);
-  return out;
-}
-
 const char* Intrinsic::str() const {
   return sem::str(type_);
 }
@@ -103,7 +98,7 @@ bool IsAtomicIntrinsic(IntrinsicType i) {
 
 Intrinsic::Intrinsic(IntrinsicType type,
                      sem::Type* return_type,
-                     const ParameterList& parameters,
+                     ParameterList parameters,
                      PipelineStageSet supported_stages,
                      bool is_deprecated)
     : Base(return_type, parameters),
@@ -111,8 +106,6 @@ Intrinsic::Intrinsic(IntrinsicType type,
       supported_stages_(supported_stages),
       is_deprecated_(is_deprecated) {}
 
-Intrinsic::Intrinsic(const Intrinsic&) = default;
-
 Intrinsic::~Intrinsic() = default;
 
 bool Intrinsic::IsCoarseDerivative() const {
@@ -155,25 +148,5 @@ bool Intrinsic::IsAtomic() const {
   return IsAtomicIntrinsic(type_);
 }
 
-bool operator==(const Intrinsic& a, const Intrinsic& b) {
-  static_assert(sizeof(Intrinsic(IntrinsicType::kNone, nullptr, ParameterList{},
-                                 PipelineStageSet{}, false)) > 0,
-                "don't forget to update the comparison below if you change the "
-                "constructor of Intrinsic!");
-
-  if (a.Type() != b.Type() || a.SupportedStages() != b.SupportedStages() ||
-      a.ReturnType() != b.ReturnType() ||
-      a.IsDeprecated() != b.IsDeprecated() ||
-      a.Parameters().size() != b.Parameters().size()) {
-    return false;
-  }
-  for (size_t i = 0; i < a.Parameters().size(); i++) {
-    if (a.Parameters()[i] != b.Parameters()[i]) {
-      return false;
-    }
-  }
-  return true;
-}
-
 }  // namespace sem
 }  // namespace tint
diff --git a/src/sem/intrinsic.h b/src/sem/intrinsic.h
index 25c0c5931d..74f332f4a8 100644
--- a/src/sem/intrinsic.h
+++ b/src/sem/intrinsic.h
@@ -88,13 +88,10 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
   /// deprecated
   Intrinsic(IntrinsicType type,
             sem::Type* return_type,
-            const ParameterList& parameters,
+            ParameterList parameters,
             PipelineStageSet supported_stages,
             bool is_deprecated);
 
-  /// Copy constructor
-  Intrinsic(const Intrinsic&);
-
   /// Destructor
   ~Intrinsic() override;
 
@@ -147,18 +144,6 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
   bool const is_deprecated_;
 };
 
-/// Emits the name of the intrinsic function type. The spelling, including case,
-/// matches the name in the WGSL spec.
-std::ostream& operator<<(std::ostream& out, IntrinsicType i);
-
-/// Equality operator for Intrinsics
-bool operator==(const Intrinsic& a, const Intrinsic& b);
-
-/// Inequality operator for Intrinsics
-static inline bool operator!=(const Intrinsic& a, const Intrinsic& b) {
-  return !(a == b);
-}
-
 }  // namespace sem
 }  // namespace tint
 
diff --git a/src/sem/intrinsic_type.cc b/src/sem/intrinsic_type.cc
index 75616caa5c..a9cac51b7d 100644
--- a/src/sem/intrinsic_type.cc
+++ b/src/sem/intrinsic_type.cc
@@ -24,6 +24,8 @@
 
 #include "src/sem/intrinsic_type.h"
 
+#include <sstream>
+
 namespace tint {
 namespace sem {
 
@@ -529,5 +531,10 @@ const char* str(IntrinsicType i) {
   return "<unknown>";
 }
 
+std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
+  out << str(i);
+  return out;
+}
+
 }  // namespace sem
 }  // namespace tint
diff --git a/src/sem/intrinsic_type.cc.tmpl b/src/sem/intrinsic_type.cc.tmpl
index e4059e3c51..f02c13f076 100644
--- a/src/sem/intrinsic_type.cc.tmpl
+++ b/src/sem/intrinsic_type.cc.tmpl
@@ -10,6 +10,8 @@ See:
 
 #include "src/sem/intrinsic_type.h"
 
+#include <sstream>
+
 namespace tint {
 namespace sem {
 
@@ -34,5 +36,10 @@ const char* str(IntrinsicType i) {
   return "<unknown>";
 }
 
+std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
+  out << str(i);
+  return out;
+}
+
 }  // namespace sem
 }  // namespace tint
diff --git a/src/sem/intrinsic_type.h b/src/sem/intrinsic_type.h
index 29e754b3d9..6e47c3e55d 100644
--- a/src/sem/intrinsic_type.h
+++ b/src/sem/intrinsic_type.h
@@ -26,6 +26,7 @@
 #define SRC_SEM_INTRINSIC_TYPE_H_
 
 #include <string>
+#include <sstream>
 
 namespace tint {
 namespace sem {
@@ -143,6 +144,10 @@ IntrinsicType ParseIntrinsicType(const std::string& name);
 /// case, matches the name in the WGSL spec.
 const char* str(IntrinsicType i);
 
+/// Emits the name of the intrinsic function type. The spelling, including case,
+/// matches the name in the WGSL spec.
+std::ostream& operator<<(std::ostream& out, IntrinsicType i);
+
 }  // namespace sem
 }  // namespace tint
 
diff --git a/src/sem/intrinsic_type.h.tmpl b/src/sem/intrinsic_type.h.tmpl
index 88826450ba..deb935c988 100644
--- a/src/sem/intrinsic_type.h.tmpl
+++ b/src/sem/intrinsic_type.h.tmpl
@@ -12,6 +12,7 @@ See:
 #define SRC_SEM_INTRINSIC_TYPE_H_
 
 #include <string>
+#include <sstream>
 
 namespace tint {
 namespace sem {
@@ -34,6 +35,10 @@ IntrinsicType ParseIntrinsicType(const std::string& name);
 /// case, matches the name in the WGSL spec.
 const char* str(IntrinsicType i);
 
+/// Emits the name of the intrinsic function type. The spelling, including case,
+/// matches the name in the WGSL spec.
+std::ostream& operator<<(std::ostream& out, IntrinsicType i);
+
 }  // namespace sem
 }  // namespace tint
 
diff --git a/src/sem/sampler_type_test.cc b/src/sem/sampler_type_test.cc
index f27ac2e0ee..9f2d5fc3a8 100644
--- a/src/sem/sampler_type_test.cc
+++ b/src/sem/sampler_type_test.cc
@@ -12,6 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
+#include "src/sem/sampler_type.h"
 #include "src/sem/test_helper.h"
 #include "src/sem/texture_type.h"
 
diff --git a/src/sem/variable.cc b/src/sem/variable.cc
index 490f890cd8..fb6f48fb51 100644
--- a/src/sem/variable.cc
+++ b/src/sem/variable.cc
@@ -20,6 +20,9 @@
 #include "src/ast/variable.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Variable);
+TINT_INSTANTIATE_TYPEINFO(tint::sem::GlobalVariable);
+TINT_INSTANTIATE_TYPEINFO(tint::sem::LocalVariable);
+TINT_INSTANTIATE_TYPEINFO(tint::sem::Parameter);
 TINT_INSTANTIATE_TYPEINFO(tint::sem::VariableUser);
 
 namespace tint {
@@ -28,26 +31,51 @@ namespace sem {
 Variable::Variable(const ast::Variable* declaration,
                    const sem::Type* type,
                    ast::StorageClass storage_class,
-                   ast::Access access,
-                   sem::BindingPoint binding_point)
+                   ast::Access access)
     : declaration_(declaration),
       type_(type),
       storage_class_(storage_class),
-      access_(access),
+      access_(access) {}
+
+Variable::~Variable() = default;
+
+LocalVariable::LocalVariable(const ast::Variable* declaration,
+                             const sem::Type* type,
+                             ast::StorageClass storage_class,
+                             ast::Access access)
+    : Base(declaration, type, storage_class, access) {}
+
+LocalVariable::~LocalVariable() = default;
+
+GlobalVariable::GlobalVariable(const ast::Variable* declaration,
+                               const sem::Type* type,
+                               ast::StorageClass storage_class,
+                               ast::Access access,
+                               sem::BindingPoint binding_point)
+    : Base(declaration, type, storage_class, access),
       binding_point_(binding_point),
       is_pipeline_constant_(false) {}
 
-Variable::Variable(const ast::Variable* declaration,
-                   const sem::Type* type,
-                   uint16_t constant_id)
-    : declaration_(declaration),
-      type_(type),
-      storage_class_(ast::StorageClass::kNone),
-      access_(ast::Access::kReadWrite),
+GlobalVariable::GlobalVariable(const ast::Variable* declaration,
+                               const sem::Type* type,
+                               uint16_t constant_id)
+    : Base(declaration,
+           type,
+           ast::StorageClass::kNone,
+           ast::Access::kReadWrite),
       is_pipeline_constant_(true),
       constant_id_(constant_id) {}
 
-Variable::~Variable() = default;
+GlobalVariable::~GlobalVariable() = default;
+
+Parameter::Parameter(const ast::Variable* declaration,
+                     const sem::Type* type,
+                     ast::StorageClass storage_class,
+                     ast::Access access,
+                     const ParameterUsage usage /* = ParameterUsage::kNone */)
+    : Base(declaration, type, storage_class, access), usage_(usage) {}
+
+Parameter::~Parameter() = default;
 
 VariableUser::VariableUser(ast::IdentifierExpression* declaration,
                            const sem::Type* type,
diff --git a/src/sem/variable.h b/src/sem/variable.h
index 356ed9b383..fc8d9beb47 100644
--- a/src/sem/variable.h
+++ b/src/sem/variable.h
@@ -21,6 +21,7 @@
 #include "src/ast/storage_class.h"
 #include "src/sem/binding_point.h"
 #include "src/sem/expression.h"
+#include "src/sem/parameter_usage.h"
 
 namespace tint {
 
@@ -36,28 +37,19 @@ namespace sem {
 class Type;
 class VariableUser;
 
-/// Variable holds the semantic information for variables.
+/// Variable is the base class for local variables, global variables and
+/// parameters.
 class Variable : public Castable<Variable, Node> {
  public:
-  /// Constructor for variables and non-overridable constants
+  /// Constructor
   /// @param declaration the AST declaration node
   /// @param type the variable type
   /// @param storage_class the variable storage class
   /// @param access the variable access control type
-  /// @param binding_point the optional resource binding point of the variable
   Variable(const ast::Variable* declaration,
            const sem::Type* type,
            ast::StorageClass storage_class,
-           ast::Access access,
-           sem::BindingPoint binding_point = {});
-
-  /// Constructor for overridable pipeline constants
-  /// @param declaration the AST declaration node
-  /// @param type the variable type
-  /// @param constant_id the pipeline constant ID
-  Variable(const ast::Variable* declaration,
-           const sem::Type* type,
-           uint16_t constant_id);
+           ast::Access access);
 
   /// Destructor
   ~Variable() override;
@@ -74,32 +66,109 @@ class Variable : public Castable<Variable, Node> {
   /// @returns the access control for the variable
   ast::Access Access() const { return access_; }
 
-  /// @returns the resource binding point for the variable
-  sem::BindingPoint BindingPoint() const { return binding_point_; }
-
   /// @returns the expressions that use the variable
   const std::vector<const VariableUser*>& Users() const { return users_; }
 
   /// @param user the user to add
   void AddUser(const VariableUser* user) { users_.emplace_back(user); }
 
-  /// @returns true if this variable is an overridable pipeline constant
-  bool IsPipelineConstant() const { return is_pipeline_constant_; }
-
-  /// @returns the pipeline constant ID associated with the variable
-  uint16_t ConstantId() const { return constant_id_; }
-
  private:
   const ast::Variable* const declaration_;
   const sem::Type* const type_;
   ast::StorageClass const storage_class_;
   ast::Access const access_;
-  sem::BindingPoint binding_point_;
   std::vector<const VariableUser*> users_;
-  const bool is_pipeline_constant_;
-  const uint16_t constant_id_ = 0;
 };
 
+/// LocalVariable is a function-scope variable
+class LocalVariable : public Castable<LocalVariable, Variable> {
+ public:
+  /// Constructor
+  /// @param declaration the AST declaration node
+  /// @param type the variable type
+  /// @param storage_class the variable storage class
+  /// @param access the variable access control type
+  LocalVariable(const ast::Variable* declaration,
+                const sem::Type* type,
+                ast::StorageClass storage_class,
+                ast::Access access);
+
+  /// Destructor
+  ~LocalVariable() override;
+};
+
+/// GlobalVariable is a module-scope variable
+class GlobalVariable : public Castable<GlobalVariable, Variable> {
+ public:
+  /// Constructor for non-overridable constants
+  /// @param declaration the AST declaration node
+  /// @param type the variable type
+  /// @param storage_class the variable storage class
+  /// @param access the variable access control type
+  /// @param binding_point the optional resource binding point of the variable
+  GlobalVariable(const ast::Variable* declaration,
+                 const sem::Type* type,
+                 ast::StorageClass storage_class,
+                 ast::Access access,
+                 sem::BindingPoint binding_point = {});
+
+  /// Constructor for overridable pipeline constants
+  /// @param declaration the AST declaration node
+  /// @param type the variable type
+  /// @param constant_id the pipeline constant ID
+  GlobalVariable(const ast::Variable* declaration,
+                 const sem::Type* type,
+                 uint16_t constant_id);
+
+  /// Destructor
+  ~GlobalVariable() override;
+
+  /// @returns the resource binding point for the variable
+  sem::BindingPoint BindingPoint() const { return binding_point_; }
+
+  /// @returns the pipeline constant ID associated with the variable
+  uint16_t ConstantId() const { return constant_id_; }
+
+  /// @returns true if this variable is an overridable pipeline constant
+  bool IsPipelineConstant() const { return is_pipeline_constant_; }
+
+ private:
+  sem::BindingPoint binding_point_;
+  bool const is_pipeline_constant_;
+  uint16_t const constant_id_ = 0;
+};
+
+/// Parameter is a function parameter
+class Parameter : public Castable<Parameter, Variable> {
+ public:
+  /// Constructor for function parameters
+  /// @param declaration the AST declaration node
+  /// @param type the variable type
+  /// @param storage_class the variable storage class
+  /// @param access the variable access control type
+  /// @param usage the semantic usage for the parameter
+  Parameter(const ast::Variable* declaration,
+            const sem::Type* type,
+            ast::StorageClass storage_class,
+            ast::Access access,
+            const ParameterUsage usage = ParameterUsage::kNone);
+
+  /// Copy constructor
+  Parameter(const Parameter&);
+
+  /// Destructor
+  ~Parameter() override;
+
+  /// @returns the semantic usage for the parameter
+  ParameterUsage Usage() const { return usage_; }
+
+ private:
+  ParameterUsage const usage_;
+};
+
+/// ParameterList is a list of Parameter
+using ParameterList = std::vector<const Parameter*>;
+
 /// VariableUser holds the semantic information for an identifier expression
 /// node that resolves to a variable.
 class VariableUser : public Castable<VariableUser, Expression> {
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index 35fbcb999f..8a0fa17a8a 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -131,7 +131,13 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
     }
 
     // Get the index to use for the buffer size array.
-    auto binding = storage_buffer_sem->Variable()->BindingPoint();
+    auto* var = tint::As<sem::GlobalVariable>(storage_buffer_sem->Variable());
+    if (!var) {
+      TINT_ICE(Transform, ctx.dst->Diagnostics())
+          << "storage buffer is not a global variable";
+      break;
+    }
+    auto binding = var->BindingPoint();
     auto idx_itr = cfg->bindpoint_to_size_index.find(binding);
     if (idx_itr == cfg->bindpoint_to_size_index.end()) {
       ctx.dst->Diagnostics().add_error(
diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index f06b3ae168..0e2b245ea6 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -682,8 +682,8 @@ struct DecomposeMemoryAccess::State {
 
       // Other parameters are copied as-is:
       for (size_t i = 1; i < intrinsic->Parameters().size(); i++) {
-        auto& param = intrinsic->Parameters()[i];
-        auto* ty = CreateASTTypeFor(ctx, param.type);
+        auto* param = intrinsic->Parameters()[i];
+        auto* ty = CreateASTTypeFor(ctx, param->Type());
         params.emplace_back(b.Param("param_" + std::to_string(i), ty));
       }
 
diff --git a/src/transform/msl.cc b/src/transform/msl.cc
index 53ead18b29..5c86a03020 100644
--- a/src/transform/msl.cc
+++ b/src/transform/msl.cc
@@ -66,10 +66,10 @@ Output Msl::Run(const Program* in, const DataMap& inputs) {
 
   // Use the SSBO binding numbers as the indices for the buffer size lookups.
   for (auto* var : in->AST().GlobalVariables()) {
-    auto* sem_var = in->Sem().Get(var);
-    if (sem_var->StorageClass() == ast::StorageClass::kStorage) {
+    auto* global = in->Sem().Get<sem::GlobalVariable>(var);
+    if (global && global->StorageClass() == ast::StorageClass::kStorage) {
       array_length_from_uniform_cfg.bindpoint_to_size_index.emplace(
-          sem_var->BindingPoint(), sem_var->BindingPoint().binding);
+          global->BindingPoint(), global->BindingPoint().binding);
     }
   }
 
diff --git a/src/transform/robustness.cc b/src/transform/robustness.cc
index 92abd7850c..1e8b94e8c9 100644
--- a/src/transform/robustness.cc
+++ b/src/transform/robustness.cc
@@ -219,7 +219,7 @@ struct Robustness::State {
 
     auto* texture_arg = expr->params()[texture_idx];
     auto* coords_arg = expr->params()[coords_idx];
-    auto* coords_ty = intrinsic->Parameters()[coords_idx].type;
+    auto* coords_ty = intrinsic->Parameters()[coords_idx]->Type();
 
     // If the level is provided, then we need to clamp this. As the level is
     // used by textureDimensions() and the texture[Load|Store]() calls, we need
diff --git a/src/transform/transform.cc b/src/transform/transform.cc
index cd788c5eae..f67c046ede 100644
--- a/src/transform/transform.cc
+++ b/src/transform/transform.cc
@@ -22,6 +22,7 @@
 #include "src/sem/block_statement.h"
 #include "src/sem/for_loop_statement.h"
 #include "src/sem/reference_type.h"
+#include "src/sem/sampler_type.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::Transform);
 TINT_INSTANTIATE_TYPEINFO(tint::transform::Data);
diff --git a/src/utils/hash.h b/src/utils/hash.h
index bd78942639..91792bf589 100644
--- a/src/utils/hash.h
+++ b/src/utils/hash.h
@@ -43,16 +43,16 @@ struct HashCombineOffset<8> {
   static constexpr inline uint64_t value() { return 0x9e3779b97f4a7c16; }
 };
 
-// When hashing sparse structures we want to iteratively build a hash value with
-// only parts of the data. HashCombine "hashes" together an existing hash and
-// hashable values.
+}  // namespace detail
+
+/// HashCombine "hashes" together an existing hash and hashable values.
 template <typename T>
 void HashCombine(size_t* hash, const T& value) {
-  constexpr size_t offset = HashCombineOffset<sizeof(size_t)>::value();
+  constexpr size_t offset = detail::HashCombineOffset<sizeof(size_t)>::value();
   *hash ^= std::hash<T>()(value) + offset + (*hash << 6) + (*hash >> 2);
 }
 
-// Helper for hashing vectors
+/// HashCombine "hashes" together an existing hash and hashable values.
 template <typename T>
 void HashCombine(size_t* hash, const std::vector<T>& vector) {
   HashCombine(hash, vector.size());
@@ -61,20 +61,19 @@ void HashCombine(size_t* hash, const std::vector<T>& vector) {
   }
 }
 
+/// HashCombine "hashes" together an existing hash and hashable values.
 template <typename T, typename... ARGS>
 void HashCombine(size_t* hash, const T& value, const ARGS&... args) {
   HashCombine(hash, value);
   HashCombine(hash, args...);
 }
 
-}  // namespace detail
-
 /// @returns a hash of the combined arguments. The returned hash is dependent on
 /// the order of the arguments.
 template <typename... ARGS>
 size_t Hash(const ARGS&... args) {
   size_t hash = 102931;  // seed with an arbitrary prime
-  detail::HashCombine(&hash, args...);
+  HashCombine(&hash, args...);
   return hash;
 }
 
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index fc9e5057e2..65b44c9fd1 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -1143,7 +1143,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
       // InterlockedExchange and discard the returned value
       {  // T result = 0;
         auto pre = line();
-        auto* value_ty = intrinsic->Parameters()[1].type;
+        auto* value_ty = intrinsic->Parameters()[1]->Type();
         if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
                              ast::Access::kUndefined, result)) {
           return false;
@@ -1278,9 +1278,9 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
-        auto* significand_ty = intrinsic->Parameters()[0].type;
+        auto* significand_ty = intrinsic->Parameters()[0]->Type();
         auto significand = params[0];
-        auto* exponent_ty = intrinsic->Parameters()[1].type;
+        auto* exponent_ty = intrinsic->Parameters()[1]->Type();
         auto exponent = params[1];
 
         std::string width;
@@ -1314,7 +1314,7 @@ bool GeneratorImpl::EmitIsNormalCall(std::ostream& out,
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
-        auto* input_ty = intrinsic->Parameters()[0].type;
+        auto* input_ty = intrinsic->Parameters()[0]->Type();
 
         std::string width;
         if (auto* vec = input_ty->As<sem::Vector>()) {
@@ -2434,12 +2434,13 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
         }
 
         if (wgsize[i].overridable_const) {
-          auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
-          if (!sem_const->IsPipelineConstant()) {
+          auto* global = builder_.Sem().Get<sem::GlobalVariable>(
+              wgsize[i].overridable_const);
+          if (!global->IsPipelineConstant()) {
             TINT_ICE(Writer, builder_.Diagnostics())
                 << "expected a pipeline-overridable constant";
           }
-          out << kSpecConstantPrefix << sem_const->ConstantId();
+          out << kSpecConstantPrefix << global->ConstantId();
         } else {
           out << std::to_string(wgsize[i].value);
         }
@@ -3187,8 +3188,9 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
   auto* sem = builder_.Sem().Get(var);
   auto* type = sem->Type();
 
-  if (sem->IsPipelineConstant()) {
-    auto const_id = sem->ConstantId();
+  auto* global = sem->As<sem::GlobalVariable>();
+  if (global && global->IsPipelineConstant()) {
+    auto const_id = global->ConstantId();
 
     line() << "#ifndef " << kSpecConstantPrefix << const_id;
 
@@ -3247,12 +3249,12 @@ bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out,
       }
       {
         ScopedParen sp(decl);
-        for (auto param : intrinsic->Parameters()) {
+        for (auto* param : intrinsic->Parameters()) {
           if (!parameter_names.empty()) {
             decl << ", ";
           }
           auto param_name = "param_" + std::to_string(parameter_names.size());
-          const auto* ty = param.type;
+          const auto* ty = param->Type();
           if (auto* ptr = ty->As<sem::Pointer>()) {
             decl << "inout ";
             ty = ptr->StoreType();
diff --git a/src/writer/hlsl/generator_impl_type_test.cc b/src/writer/hlsl/generator_impl_type_test.cc
index feedf4e841..7db8426653 100644
--- a/src/writer/hlsl/generator_impl_type_test.cc
+++ b/src/writer/hlsl/generator_impl_type_test.cc
@@ -19,6 +19,7 @@
 #include "src/sem/depth_texture_type.h"
 #include "src/sem/multisampled_texture_type.h"
 #include "src/sem/sampled_texture_type.h"
+#include "src/sem/sampler_type.h"
 #include "src/sem/storage_texture_type.h"
 #include "src/writer/hlsl/test_helper.h"
 
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 39292bd58f..c22b146a33 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -2323,9 +2323,9 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
     out << " " << program_->Symbols().NameFor(var->symbol());
   }
 
-  auto* sem_var = program_->Sem().Get(var);
-  if (sem_var->IsPipelineConstant()) {
-    out << " [[function_constant(" << sem_var->ConstantId() << ")]]";
+  auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
+  if (global && global->IsPipelineConstant()) {
+    out << " [[function_constant(" << global->ConstantId() << ")]]";
   } else if (var->constructor() != nullptr) {
     out << " = ";
     if (!EmitExpression(out, var->constructor())) {
diff --git a/src/writer/msl/generator_impl_type_test.cc b/src/writer/msl/generator_impl_type_test.cc
index fa55b75814..646ae7a349 100644
--- a/src/writer/msl/generator_impl_type_test.cc
+++ b/src/writer/msl/generator_impl_type_test.cc
@@ -20,6 +20,7 @@
 #include "src/sem/depth_texture_type.h"
 #include "src/sem/multisampled_texture_type.h"
 #include "src/sem/sampled_texture_type.h"
+#include "src/sem/sampler_type.h"
 #include "src/sem/storage_texture_type.h"
 #include "src/writer/msl/test_helper.h"
 
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index d9bf3633c5..f8e8ead53d 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -492,7 +492,8 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
         auto constant = ScalarConstant::U32(wgsize[i].value);
         if (wgsize[i].overridable_const) {
           // Make the constant specializable.
-          auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
+          auto* sem_const = builder_.Sem().Get<sem::GlobalVariable>(
+              wgsize[i].overridable_const);
           if (!sem_const->IsPipelineConstant()) {
             TINT_ICE(Writer, builder_.Diagnostics())
                 << "expected a pipeline-overridable constant";
@@ -1635,10 +1636,10 @@ uint32_t Builder::GenerateLiteralIfNeeded(ast::Variable* var,
                                           ast::Literal* lit) {
   ScalarConstant constant;
 
-  auto* sem_var = builder_.Sem().Get(var);
-  if (sem_var && sem_var->IsPipelineConstant()) {
+  auto* global = builder_.Sem().Get<sem::GlobalVariable>(var);
+  if (global && global->IsPipelineConstant()) {
     constant.is_spec_op = true;
-    constant.constant_id = sem_var->ConstantId();
+    constant.constant_id = global->ConstantId();
   }
 
   if (auto* l = lit->As<ast::BoolLiteral>()) {
@@ -2295,13 +2296,13 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
   // and loads it if necessary. Returns 0 on error.
   auto get_param_as_value_id = [&](size_t i) -> uint32_t {
     auto* arg = call->params()[i];
-    auto& param = intrinsic->Parameters()[i];
+    auto* param = intrinsic->Parameters()[i];
     auto val_id = GenerateExpression(arg);
     if (val_id == 0) {
       return 0;
     }
 
-    if (!param.type->Is<sem::Pointer>()) {
+    if (!param->Type()->Is<sem::Pointer>()) {
       val_id = GenerateLoadIfNeeded(TypeOf(arg), val_id);
     }
     return val_id;
@@ -2527,7 +2528,8 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
       // splat the condition into a vector of the same size.
       // TODO(jrprice): If we're targeting SPIR-V 1.4, we don't need to do this.
       auto* result_vector_type = intrinsic->ReturnType()->As<sem::Vector>();
-      if (result_vector_type && intrinsic->Parameters()[2].type->is_scalar()) {
+      if (result_vector_type &&
+          intrinsic->Parameters()[2]->Type()->is_scalar()) {
         sem::Bool bool_type;
         sem::Vector bool_vec_type(&bool_type, result_vector_type->size());
         if (!GenerateTypeIfNeeded(&bool_vec_type)) {
@@ -3036,14 +3038,15 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
                                       Operand result_type,
                                       Operand result_id) {
   auto is_value_signed = [&] {
-    return intrinsic->Parameters()[1].type->Is<sem::I32>();
+    return intrinsic->Parameters()[1]->Type()->Is<sem::I32>();
   };
 
   auto storage_class =
-      intrinsic->Parameters()[0].type->As<sem::Pointer>()->StorageClass();
+      intrinsic->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass();
 
   uint32_t memory_id = 0;
-  switch (intrinsic->Parameters()[0].type->As<sem::Pointer>()->StorageClass()) {
+  switch (
+      intrinsic->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass()) {
     case ast::StorageClass::kWorkgroup:
       memory_id = GenerateConstantIfNeeded(
           ScalarConstant::U32(static_cast<uint32_t>(spv::Scope::Workgroup)));
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index a81da1e47a..c33db7c8b4 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -35,6 +35,7 @@
 #include "src/ast/variable_decl_statement.h"
 #include "src/program_builder.h"
 #include "src/scope_stack.h"
+#include "src/sem/intrinsic.h"
 #include "src/sem/storage_texture_type.h"
 #include "src/writer/spirv/function.h"
 #include "src/writer/spirv/scalar_constant.h"