diff --git a/BUILD.gn b/BUILD.gn
index 8e15afc8b5..12578bb58e 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -385,6 +385,7 @@ source_set("libtint_core_src") {
     "src/semantic/sem_function.cc",
     "src/semantic/sem_info.cc",
     "src/semantic/sem_node.cc",
+    "src/semantic/sem_variable.cc",
     "src/semantic/type_mappings.h",
     "src/source.cc",
     "src/source.h",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 56a1efb9da..f79ae1a891 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -199,6 +199,7 @@ set(TINT_LIB_SRCS
   semantic/sem_function.cc
   semantic/sem_info.cc
   semantic/sem_node.cc
+  semantic/sem_variable.cc
   semantic/type_mappings.h
   source.cc
   source.h
diff --git a/src/ast/variable.cc b/src/ast/variable.cc
index cc0e676491..40056a4fb6 100644
--- a/src/ast/variable.cc
+++ b/src/ast/variable.cc
@@ -19,6 +19,7 @@
 #include "src/ast/constant_id_decoration.h"
 #include "src/clone_context.h"
 #include "src/program_builder.h"
+#include "src/semantic/variable.h"
 
 TINT_INSTANTIATE_CLASS_ID(tint::ast::Variable);
 
@@ -38,7 +39,7 @@ Variable::Variable(const Source& source,
       is_const_(is_const),
       constructor_(constructor),
       decorations_(std::move(decorations)),
-      storage_class_(sc) {}
+      declared_storage_class_(sc) {}
 
 Variable::Variable(Variable&&) = default;
 
@@ -91,10 +92,10 @@ uint32_t Variable::constant_id() const {
 }
 
 Variable* Variable::Clone(CloneContext* ctx) const {
-  return ctx->dst->create<Variable>(ctx->Clone(source()), ctx->Clone(symbol_),
-                                    storage_class(), ctx->Clone(type()),
-                                    is_const_, ctx->Clone(constructor()),
-                                    ctx->Clone(decorations_));
+  return ctx->dst->create<Variable>(
+      ctx->Clone(source()), ctx->Clone(symbol_), declared_storage_class(),
+      ctx->Clone(type()), is_const_, ctx->Clone(constructor()),
+      ctx->Clone(decorations_));
 }
 
 bool Variable::IsValid() const {
@@ -110,13 +111,15 @@ bool Variable::IsValid() const {
   return true;
 }
 
-void Variable::info_to_str(const semantic::Info&,
+void Variable::info_to_str(const semantic::Info& sem,
                            std::ostream& out,
                            size_t indent) const {
+  auto* var_sem = sem.Get(this);
   make_indent(out, indent);
   out << symbol_.to_str() << std::endl;
   make_indent(out, indent);
-  out << storage_class_ << std::endl;
+  out << (var_sem ? var_sem->StorageClass() : declared_storage_class())
+      << std::endl;
   make_indent(out, indent);
   out << type_->type_name() << std::endl;
 }
diff --git a/src/ast/variable.h b/src/ast/variable.h
index fb337b1f56..2d18b14ebf 100644
--- a/src/ast/variable.h
+++ b/src/ast/variable.h
@@ -84,7 +84,7 @@ class Variable : public Castable<Variable, Node> {
   /// Create a variable
   /// @param source the variable source
   /// @param sym the variable symbol
-  /// @param sc the variable storage class
+  /// @param sc the declared storage class
   /// @param type the value type
   /// @param is_const true if the variable is const
   /// @param constructor the constructor expression
@@ -107,12 +107,10 @@ class Variable : public Castable<Variable, Node> {
   /// @returns the variable's type.
   type::Type* type() const { return type_; }
 
-  /// Sets the storage class
-  /// @param sc the storage class
-  void set_storage_class(StorageClass sc) { storage_class_ = sc; }
-  /// @returns the storage class
-  StorageClass storage_class() const { return storage_class_; }
-
+  /// @returns the declared storage class
+  StorageClass declared_storage_class() const {
+    return declared_storage_class_;
+  }
   /// @returns the constructor expression or nullptr if none set
   Expression* constructor() const { return constructor_; }
   /// @returns true if the variable has an constructor
@@ -126,7 +124,7 @@ class Variable : public Castable<Variable, Node> {
 
   /// @returns true if the decorations include a LocationDecoration
   bool HasLocationDecoration() const;
-  /// @returns true if the deocrations include a BuiltinDecoration
+  /// @returns true if the decorations include a BuiltinDecoration
   bool HasBuiltinDecoration() const;
   /// @returns true if the decorations include a ConstantIdDecoration
   bool HasConstantIdDecoration() const;
@@ -182,8 +180,7 @@ class Variable : public Castable<Variable, Node> {
   bool const is_const_;
   Expression* const constructor_;
   VariableDecorationList const decorations_;
-
-  StorageClass storage_class_ = StorageClass::kNone;  // Semantic info
+  StorageClass const declared_storage_class_;
 };
 
 /// A list of variables
diff --git a/src/ast/variable_test.cc b/src/ast/variable_test.cc
index f700381d57..9a9f6be27b 100644
--- a/src/ast/variable_test.cc
+++ b/src/ast/variable_test.cc
@@ -30,7 +30,7 @@ TEST_F(VariableTest, Creation) {
   auto* v = Var("my_var", StorageClass::kFunction, ty.i32());
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->storage_class(), StorageClass::kFunction);
+  EXPECT_EQ(v->declared_storage_class(), StorageClass::kFunction);
   EXPECT_EQ(v->type(), ty.i32());
   EXPECT_EQ(v->source().range.begin.line, 0u);
   EXPECT_EQ(v->source().range.begin.column, 0u);
@@ -44,7 +44,7 @@ TEST_F(VariableTest, CreationWithSource) {
       "i", StorageClass::kPrivate, ty.f32(), nullptr, VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->storage_class(), StorageClass::kPrivate);
+  EXPECT_EQ(v->declared_storage_class(), StorageClass::kPrivate);
   EXPECT_EQ(v->type(), ty.f32());
   EXPECT_EQ(v->source().range.begin.line, 27u);
   EXPECT_EQ(v->source().range.begin.column, 4u);
@@ -59,7 +59,7 @@ TEST_F(VariableTest, CreationEmpty) {
       VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->storage_class(), StorageClass::kWorkgroup);
+  EXPECT_EQ(v->declared_storage_class(), StorageClass::kWorkgroup);
   EXPECT_EQ(v->type(), ty.i32());
   EXPECT_EQ(v->source().range.begin.line, 27u);
   EXPECT_EQ(v->source().range.begin.column, 4u);
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index ccfe471bc9..ca4bb9767f 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -30,6 +30,7 @@
 #include "src/ast/variable.h"
 #include "src/program.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/array_type.h"
 #include "src/type/f32_type.h"
@@ -66,14 +67,16 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
              entry_point.workgroup_size_z) = func->workgroup_size();
 
     for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) {
-      auto name = program_->Symbols().NameFor(var->symbol());
-      if (var->HasBuiltinDecoration()) {
+      auto* decl = var->Declaration();
+
+      auto name = program_->Symbols().NameFor(decl->symbol());
+      if (decl->HasBuiltinDecoration()) {
         continue;
       }
 
       StageVariable stage_variable;
       stage_variable.name = name;
-      auto* location_decoration = var->GetLocationDecoration();
+      auto* location_decoration = decl->GetLocationDecoration();
       if (location_decoration) {
         stage_variable.has_location_decoration = true;
         stage_variable.location_decoration = location_decoration->value();
@@ -81,9 +84,9 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
         stage_variable.has_location_decoration = false;
       }
 
-      if (var->storage_class() == ast::StorageClass::kInput) {
+      if (var->StorageClass() == ast::StorageClass::kInput) {
         entry_point.input_variables.push_back(stage_variable);
-      } else if (var->storage_class() == ast::StorageClass::kOutput) {
+      } else if (var->StorageClass() == ast::StorageClass::kOutput) {
         entry_point.output_variables.push_back(stage_variable);
       }
     }
@@ -188,14 +191,14 @@ std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
 
   auto* func_sem = program_->Sem().Get(func);
   for (auto& ruv : func_sem->ReferencedUniformVariables()) {
-    ResourceBinding entry;
-    ast::Variable* var = nullptr;
-    semantic::Function::BindingInfo binding_info;
-    std::tie(var, binding_info) = ruv;
-    if (!var->type()->Is<type::AccessControl>()) {
+    auto* var = ruv.first;
+    auto* decl = var->Declaration();
+    auto binding_info = ruv.second;
+
+    if (!decl->type()->Is<type::AccessControl>()) {
       continue;
     }
-    auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+    auto* unwrapped_type = decl->type()->UnwrapIfNeeded();
 
     auto* str = unwrapped_type->As<type::Struct>();
     if (str == nullptr) {
@@ -206,10 +209,11 @@ std::vector<ResourceBinding> Inspector::GetUniformBufferResourceBindings(
       continue;
     }
 
+    ResourceBinding entry;
     entry.bind_group = binding_info.group->value();
     entry.binding = binding_info.binding->value();
     entry.min_buffer_binding_size =
-        var->type()->MinBufferBindingSize(type::MemoryLayout::kUniformBuffer);
+        decl->type()->MinBufferBindingSize(type::MemoryLayout::kUniformBuffer);
 
     result.push_back(entry);
   }
@@ -239,11 +243,9 @@ std::vector<ResourceBinding> Inspector::GetSamplerResourceBindings(
 
   auto* func_sem = program_->Sem().Get(func);
   for (auto& rs : func_sem->ReferencedSamplerVariables()) {
-    ResourceBinding entry;
-    ast::Variable* var = nullptr;
-    semantic::Function::BindingInfo binding_info;
-    std::tie(var, binding_info) = rs;
+    auto binding_info = rs.second;
 
+    ResourceBinding entry;
     entry.bind_group = binding_info.group->value();
     entry.binding = binding_info.binding->value();
 
@@ -264,11 +266,9 @@ std::vector<ResourceBinding> Inspector::GetComparisonSamplerResourceBindings(
 
   auto* func_sem = program_->Sem().Get(func);
   for (auto& rcs : func_sem->ReferencedComparisonSamplerVariables()) {
-    ResourceBinding entry;
-    ast::Variable* var = nullptr;
-    semantic::Function::BindingInfo binding_info;
-    std::tie(var, binding_info) = rcs;
+    auto binding_info = rcs.second;
 
+    ResourceBinding entry;
     entry.bind_group = binding_info.group->value();
     entry.binding = binding_info.binding->value();
 
@@ -314,12 +314,11 @@ std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
   auto* func_sem = program_->Sem().Get(func);
   std::vector<ResourceBinding> result;
   for (auto& rsv : func_sem->ReferencedStoragebufferVariables()) {
-    ResourceBinding entry;
-    ast::Variable* var = nullptr;
-    semantic::Function::BindingInfo binding_info;
-    std::tie(var, binding_info) = rsv;
+    auto* var = rsv.first;
+    auto* decl = var->Declaration();
+    auto binding_info = rsv.second;
 
-    auto* ac_type = var->type()->As<type::AccessControl>();
+    auto* ac_type = decl->type()->As<type::AccessControl>();
     if (ac_type == nullptr) {
       continue;
     }
@@ -328,14 +327,15 @@ std::vector<ResourceBinding> Inspector::GetStorageBufferResourceBindingsImpl(
       continue;
     }
 
-    if (!var->type()->UnwrapIfNeeded()->Is<type::Struct>()) {
+    if (!decl->type()->UnwrapIfNeeded()->Is<type::Struct>()) {
       continue;
     }
 
+    ResourceBinding entry;
     entry.bind_group = binding_info.group->value();
     entry.binding = binding_info.binding->value();
     entry.min_buffer_binding_size =
-        var->type()->MinBufferBindingSize(type::MemoryLayout::kStorageBuffer);
+        decl->type()->MinBufferBindingSize(type::MemoryLayout::kStorageBuffer);
 
     result.push_back(entry);
   }
@@ -357,15 +357,15 @@ std::vector<ResourceBinding> Inspector::GetSampledTextureResourceBindingsImpl(
       multisampled_only ? func_sem->ReferencedMultisampledTextureVariables()
                         : func_sem->ReferencedSampledTextureVariables();
   for (auto& ref : referenced_variables) {
-    ResourceBinding entry;
-    ast::Variable* var = nullptr;
-    semantic::Function::BindingInfo binding_info;
-    std::tie(var, binding_info) = ref;
+    auto* var = ref.first;
+    auto* decl = var->Declaration();
+    auto binding_info = ref.second;
 
+    ResourceBinding entry;
     entry.bind_group = binding_info.group->value();
     entry.binding = binding_info.binding->value();
 
-    auto* texture_type = var->type()->UnwrapIfNeeded()->As<type::Texture>();
+    auto* texture_type = decl->type()->UnwrapIfNeeded()->As<type::Texture>();
     switch (texture_type->dim()) {
       case type::TextureDimension::k1d:
         entry.dim = ResourceBinding::TextureDimension::k1d;
diff --git a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
index 9c83e4221a..625abb59b4 100644
--- a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
@@ -38,7 +38,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithoutConstructor) {
 
   EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
   EXPECT_TRUE(e->type()->Is<type::F32>());
-  EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
+  EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput);
 
   EXPECT_EQ(e->source().range.begin.line, 1u);
   EXPECT_EQ(e->source().range.begin.column, 10u);
@@ -61,7 +61,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithConstructor) {
 
   EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
   EXPECT_TRUE(e->type()->Is<type::F32>());
-  EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
+  EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput);
 
   EXPECT_EQ(e->source().range.begin.line, 1u);
   EXPECT_EQ(e->source().range.begin.column, 10u);
@@ -87,7 +87,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration) {
   EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
   ASSERT_NE(e->type(), nullptr);
   EXPECT_TRUE(e->type()->Is<type::F32>());
-  EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
+  EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput);
 
   EXPECT_EQ(e->source().range.begin.line, 1u);
   EXPECT_EQ(e->source().range.begin.column, 35u);
@@ -117,7 +117,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration_MulitpleGroups) {
   EXPECT_EQ(e->symbol(), p->builder().Symbols().Get("a"));
   ASSERT_NE(e->type(), nullptr);
   EXPECT_TRUE(e->type()->Is<type::F32>());
-  EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
+  EXPECT_EQ(e->declared_storage_class(), ast::StorageClass::kOutput);
 
   EXPECT_EQ(e->source().range.begin.line, 1u);
   EXPECT_EQ(e->source().range.begin.column, 38u);
diff --git a/src/semantic/function.h b/src/semantic/function.h
index 0c751138b9..70708d986c 100644
--- a/src/semantic/function.h
+++ b/src/semantic/function.h
@@ -28,7 +28,6 @@ namespace tint {
 namespace ast {
 class BindingDecoration;
 class GroupDecoration;
-class Variable;
 class LocationDecoration;
 class BuiltinDecoration;
 }  // namespace ast
@@ -38,6 +37,8 @@ class Type;
 
 namespace semantic {
 
+class Variable;
+
 /// Function holds the semantic information for function nodes.
 class Function : public Castable<Function, Node> {
  public:
@@ -54,8 +55,8 @@ class Function : public Castable<Function, Node> {
   /// @param local_referenced_module_vars the locally referenced module
   /// variables
   /// @param ancestor_entry_points the ancestor entry points
-  explicit Function(std::vector<ast::Variable*> referenced_module_vars,
-                    std::vector<ast::Variable*> local_referenced_module_vars,
+  explicit Function(std::vector<const Variable*> referenced_module_vars,
+                    std::vector<const Variable*> local_referenced_module_vars,
                     std::vector<Symbol> ancestor_entry_points);
 
   /// Destructor
@@ -64,11 +65,11 @@ class Function : public Castable<Function, Node> {
   /// 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
-  const std::vector<ast::Variable*>& ReferencedModuleVariables() const {
+  const std::vector<const Variable*>& ReferencedModuleVariables() const {
     return referenced_module_vars_;
   }
   /// @returns the locally referenced module variables
-  const std::vector<ast::Variable*>& LocalReferencedModuleVariables() const {
+  const std::vector<const Variable*>& LocalReferencedModuleVariables() const {
     return local_referenced_module_vars_;
   }
   /// @returns the ancestor entry points
@@ -77,53 +78,53 @@ class Function : public Castable<Function, Node> {
   }
   /// Retrieves any referenced location variables
   /// @returns the <variable, decoration> pair.
-  const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
+  const std::vector<std::pair<const Variable*, ast::LocationDecoration*>>
   ReferencedLocationVariables() const;
 
   /// Retrieves any referenced builtin variables
   /// @returns the <variable, decoration> pair.
-  const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+  const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
   ReferencedBuiltinVariables() const;
 
   /// Retrieves any referenced uniform variables. Note, the variables must be
   /// decorated with both binding and group decorations.
   /// @returns the referenced uniforms
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedUniformVariables() const;
 
   /// Retrieves any referenced storagebuffer variables. Note, the variables
   /// must be decorated with both binding and group decorations.
   /// @returns the referenced storagebuffers
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedStoragebufferVariables() const;
 
   /// Retrieves any referenced regular Sampler variables. Note, the
   /// variables must be decorated with both binding and group decorations.
   /// @returns the referenced storagebuffers
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedSamplerVariables() const;
 
   /// Retrieves any referenced comparison Sampler variables. Note, the
   /// variables must be decorated with both binding and group decorations.
   /// @returns the referenced storagebuffers
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedComparisonSamplerVariables() const;
 
   /// Retrieves any referenced sampled textures variables. Note, the
   /// variables must be decorated with both binding and group decorations.
   /// @returns the referenced sampled textures
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedSampledTextureVariables() const;
 
   /// Retrieves any referenced multisampled textures variables. Note, the
   /// variables must be decorated with both binding and group decorations.
   /// @returns the referenced sampled textures
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedMultisampledTextureVariables() const;
 
   /// Retrieves any locally referenced builtin variables
   /// @returns the <variable, decoration> pairs.
-  const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+  const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
   LocalReferencedBuiltinVariables() const;
 
   /// Checks if the given entry point is an ancestor
@@ -132,13 +133,13 @@ class Function : public Castable<Function, Node> {
   bool HasAncestorEntryPoint(Symbol sym) const;
 
  private:
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedSamplerVariablesImpl(type::SamplerKind kind) const;
-  const std::vector<std::pair<ast::Variable*, BindingInfo>>
+  const std::vector<std::pair<const Variable*, BindingInfo>>
   ReferencedSampledTextureVariablesImpl(bool multisampled) const;
 
-  std::vector<ast::Variable*> const referenced_module_vars_;
-  std::vector<ast::Variable*> const local_referenced_module_vars_;
+  std::vector<const Variable*> const referenced_module_vars_;
+  std::vector<const Variable*> const local_referenced_module_vars_;
   std::vector<Symbol> const ancestor_entry_points_;
 };
 
diff --git a/src/semantic/sem_function.cc b/src/semantic/sem_function.cc
index 561369702f..26768ec9fc 100644
--- a/src/semantic/sem_function.cc
+++ b/src/semantic/sem_function.cc
@@ -20,6 +20,7 @@
 #include "src/ast/location_decoration.h"
 #include "src/ast/variable.h"
 #include "src/ast/variable_decoration.h"
+#include "src/semantic/variable.h"
 #include "src/type/multisampled_texture_type.h"
 #include "src/type/sampled_texture_type.h"
 #include "src/type/texture_type.h"
@@ -29,8 +30,8 @@ TINT_INSTANTIATE_CLASS_ID(tint::semantic::Function);
 namespace tint {
 namespace semantic {
 
-Function::Function(std::vector<ast::Variable*> referenced_module_vars,
-                   std::vector<ast::Variable*> local_referenced_module_vars,
+Function::Function(std::vector<const Variable*> referenced_module_vars,
+                   std::vector<const Variable*> local_referenced_module_vars,
                    std::vector<Symbol> ancestor_entry_points)
     : referenced_module_vars_(std::move(referenced_module_vars)),
       local_referenced_module_vars_(std::move(local_referenced_module_vars)),
@@ -38,12 +39,12 @@ Function::Function(std::vector<ast::Variable*> referenced_module_vars,
 
 Function::~Function() = default;
 
-const std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>>
+const std::vector<std::pair<const Variable*, ast::LocationDecoration*>>
 Function::ReferencedLocationVariables() const {
-  std::vector<std::pair<ast::Variable*, ast::LocationDecoration*>> ret;
+  std::vector<std::pair<const Variable*, ast::LocationDecoration*>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* location = deco->As<ast::LocationDecoration>()) {
         ret.push_back({var, location});
         break;
@@ -53,18 +54,18 @@ Function::ReferencedLocationVariables() const {
   return ret;
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedUniformVariables() const {
-  std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+  std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    if (var->storage_class() != ast::StorageClass::kUniform) {
+    if (var->StorageClass() != ast::StorageClass::kUniform) {
       continue;
     }
 
     ast::BindingDecoration* binding = nullptr;
     ast::GroupDecoration* group = nullptr;
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* b = deco->As<ast::BindingDecoration>()) {
         binding = b;
       } else if (auto* g = deco->As<ast::GroupDecoration>()) {
@@ -80,18 +81,18 @@ Function::ReferencedUniformVariables() const {
   return ret;
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedStoragebufferVariables() const {
-  std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+  std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    if (var->storage_class() != ast::StorageClass::kStorage) {
+    if (var->StorageClass() != ast::StorageClass::kStorage) {
       continue;
     }
 
     ast::BindingDecoration* binding = nullptr;
     ast::GroupDecoration* group = nullptr;
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* b = deco->As<ast::BindingDecoration>()) {
         binding = b;
       } else if (auto* s = deco->As<ast::GroupDecoration>()) {
@@ -107,12 +108,12 @@ Function::ReferencedStoragebufferVariables() const {
   return ret;
 }
 
-const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
 Function::ReferencedBuiltinVariables() const {
-  std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
+  std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
         ret.push_back({var, builtin});
         break;
@@ -122,32 +123,32 @@ Function::ReferencedBuiltinVariables() const {
   return ret;
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedSamplerVariables() const {
   return ReferencedSamplerVariablesImpl(type::SamplerKind::kSampler);
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedComparisonSamplerVariables() const {
   return ReferencedSamplerVariablesImpl(type::SamplerKind::kComparisonSampler);
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedSampledTextureVariables() const {
   return ReferencedSampledTextureVariablesImpl(false);
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedMultisampledTextureVariables() const {
   return ReferencedSampledTextureVariablesImpl(true);
 }
 
-const std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>>
+const std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>>
 Function::LocalReferencedBuiltinVariables() const {
-  std::vector<std::pair<ast::Variable*, ast::BuiltinDecoration*>> ret;
+  std::vector<std::pair<const Variable*, ast::BuiltinDecoration*>> ret;
 
   for (auto* var : LocalReferencedModuleVariables()) {
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
         ret.push_back({var, builtin});
         break;
@@ -166,12 +167,12 @@ bool Function::HasAncestorEntryPoint(Symbol symbol) const {
   return false;
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
-  std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+  std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+    auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded();
     auto* sampler = unwrapped_type->As<type::Sampler>();
     if (sampler == nullptr || sampler->kind() != kind) {
       continue;
@@ -179,7 +180,7 @@ Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
 
     ast::BindingDecoration* binding = nullptr;
     ast::GroupDecoration* group = nullptr;
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* b = deco->As<ast::BindingDecoration>()) {
         binding = b;
       }
@@ -196,12 +197,12 @@ Function::ReferencedSamplerVariablesImpl(type::SamplerKind kind) const {
   return ret;
 }
 
-const std::vector<std::pair<ast::Variable*, Function::BindingInfo>>
+const std::vector<std::pair<const Variable*, Function::BindingInfo>>
 Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const {
-  std::vector<std::pair<ast::Variable*, Function::BindingInfo>> ret;
+  std::vector<std::pair<const Variable*, Function::BindingInfo>> ret;
 
   for (auto* var : ReferencedModuleVariables()) {
-    auto* unwrapped_type = var->type()->UnwrapIfNeeded();
+    auto* unwrapped_type = var->Declaration()->type()->UnwrapIfNeeded();
     auto* texture = unwrapped_type->As<type::Texture>();
     if (texture == nullptr) {
       continue;
@@ -216,7 +217,7 @@ Function::ReferencedSampledTextureVariablesImpl(bool multisampled) const {
 
     ast::BindingDecoration* binding = nullptr;
     ast::GroupDecoration* group = nullptr;
-    for (auto* deco : var->decorations()) {
+    for (auto* deco : var->Declaration()->decorations()) {
       if (auto* b = deco->As<ast::BindingDecoration>()) {
         binding = b;
       } else if (auto* s = deco->As<ast::GroupDecoration>()) {
diff --git a/src/semantic/sem_variable.cc b/src/semantic/sem_variable.cc
new file mode 100644
index 0000000000..33e84a1d82
--- /dev/null
+++ b/src/semantic/sem_variable.cc
@@ -0,0 +1,28 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/semantic/variable.h"
+
+TINT_INSTANTIATE_CLASS_ID(tint::semantic::Variable);
+
+namespace tint {
+namespace semantic {
+
+Variable::Variable(ast::Variable* declaration, ast::StorageClass storage_class)
+    : declaration_(declaration), storage_class_(storage_class) {}
+
+Variable::~Variable() = default;
+
+}  // namespace semantic
+}  // namespace tint
diff --git a/src/semantic/type_mappings.h b/src/semantic/type_mappings.h
index ac099a8c59..58b91c5428 100644
--- a/src/semantic/type_mappings.h
+++ b/src/semantic/type_mappings.h
@@ -24,6 +24,7 @@ namespace ast {
 
 class Expression;
 class Function;
+class Variable;
 
 }  // namespace ast
 
@@ -31,6 +32,7 @@ namespace semantic {
 
 class Expression;
 class Function;
+class Variable;
 
 /// TypeMappings is a struct that holds dummy `operator()` methods that's used
 /// by SemanticNodeTypeFor to map AST node types to their corresponding semantic
@@ -41,6 +43,7 @@ struct TypeMappings {
   //! @cond Doxygen_Suppress
   semantic::Expression* operator()(ast::Expression*);
   semantic::Function* operator()(ast::Function*);
+  semantic::Variable* operator()(ast::Variable*);
   //! @endcond
 };
 
diff --git a/src/semantic/variable.h b/src/semantic/variable.h
new file mode 100644
index 0000000000..6f15dc7506
--- /dev/null
+++ b/src/semantic/variable.h
@@ -0,0 +1,64 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0(the "License");
+
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_SEMANTIC_VARIABLE_H_
+#define SRC_SEMANTIC_VARIABLE_H_
+
+#include <utility>
+#include <vector>
+
+#include "src/ast/storage_class.h"
+#include "src/semantic/node.h"
+#include "src/type/sampler_type.h"
+
+namespace tint {
+
+// Forward declarations
+namespace ast {
+class Variable;
+}  // namespace ast
+namespace type {
+class Type;
+}  // namespace type
+
+namespace semantic {
+
+/// Variable holds the semantic information for variables.
+class Variable : public Castable<Variable, Node> {
+ public:
+  /// Constructor
+  /// @param declaration the AST declaration node
+  /// @param storage_class the variable storage class
+  explicit Variable(ast::Variable* declaration,
+                    ast::StorageClass storage_class);
+
+  /// Destructor
+  ~Variable() override;
+
+  /// @returns the AST declaration node
+  ast::Variable* Declaration() const { return declaration_; }
+
+  /// @returns the storage class for the variable
+  ast::StorageClass StorageClass() const { return storage_class_; }
+
+ private:
+  ast::Variable* const declaration_;
+  ast::StorageClass const storage_class_;
+};
+
+}  // namespace semantic
+}  // namespace tint
+
+#endif  // SRC_SEMANTIC_VARIABLE_H_
diff --git a/src/transform/emit_vertex_point_size.cc b/src/transform/emit_vertex_point_size.cc
index 319d14434a..748021cd49 100644
--- a/src/transform/emit_vertex_point_size.cc
+++ b/src/transform/emit_vertex_point_size.cc
@@ -58,7 +58,6 @@ Transform::Output EmitVertexPointSize::Run(const Program* in) {
       false,                                  // is_const
       nullptr,                                // constructor
       ast::VariableDecorationList{
-          // decorations
           out.create<ast::BuiltinDecoration>(Source{},
                                              ast::Builtin::kPointSize),
       });
diff --git a/src/transform/first_index_offset.cc b/src/transform/first_index_offset.cc
index fec1068960..4ed52ab437 100644
--- a/src/transform/first_index_offset.cc
+++ b/src/transform/first_index_offset.cc
@@ -70,7 +70,7 @@ ast::Variable* clone_variable_with_new_name(CloneContext* ctx,
   return ctx->dst->create<ast::Variable>(
       ctx->Clone(in->source()),                // source
       ctx->dst->Symbols().Register(new_name),  // symbol
-      in->storage_class(),                     // storage_class
+      in->declared_storage_class(),            // declared_storage_class
       ctx->Clone(in->type()),                  // type
       in->is_const(),                          // is_const
       ctx->Clone(in->constructor()),           // constructor
@@ -226,7 +226,7 @@ ast::Variable* FirstIndexOffset::AddUniformBuffer(ProgramBuilder* dst) {
       ast::VariableDecorationList{
           dst->create<ast::BindingDecoration>(Source{}, binding_),
           dst->create<ast::GroupDecoration>(Source{}, group_),
-      });  // decorations
+      });
 
   dst->AST().AddGlobalVariable(idx_var);
 
diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc
index 39585bea7f..f052bd44aa 100644
--- a/src/transform/vertex_pulling.cc
+++ b/src/transform/vertex_pulling.cc
@@ -35,6 +35,7 @@
 #include "src/clone_context.h"
 #include "src/program.h"
 #include "src/program_builder.h"
+#include "src/semantic/variable.h"
 #include "src/type/array_type.h"
 #include "src/type/f32_type.h"
 #include "src/type/i32_type.h"
@@ -150,7 +151,8 @@ void VertexPulling::State::FindOrInsertVertexIndexIfUsed() {
 
   // Look for an existing vertex index builtin
   for (auto* v : ctx.src->AST().GlobalVariables()) {
-    if (v->storage_class() != ast::StorageClass::kInput) {
+    auto* sem = ctx.src->Sem().Get(v);
+    if (sem->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
 
@@ -196,7 +198,8 @@ void VertexPulling::State::FindOrInsertInstanceIndexIfUsed() {
 
   // Look for an existing instance index builtin
   for (auto* v : ctx.src->AST().GlobalVariables()) {
-    if (v->storage_class() != ast::StorageClass::kInput) {
+    auto* sem = ctx.src->Sem().Get(v);
+    if (sem->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
 
@@ -229,7 +232,8 @@ void VertexPulling::State::FindOrInsertInstanceIndexIfUsed() {
 
 void VertexPulling::State::ConvertVertexInputVariablesToPrivate() {
   for (auto* v : ctx.src->AST().GlobalVariables()) {
-    if (v->storage_class() != ast::StorageClass::kInput) {
+    auto* sem = ctx.src->Sem().Get(v);
+    if (sem->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
 
diff --git a/src/type_determiner.cc b/src/type_determiner.cc
index 61723843ab..e75bdc8f49 100644
--- a/src/type_determiner.cc
+++ b/src/type_determiner.cc
@@ -45,6 +45,7 @@
 #include "src/program_builder.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/array_type.h"
 #include "src/type/bool_type.h"
 #include "src/type/depth_texture_type.h"
@@ -88,13 +89,13 @@ void TypeDeterminer::set_error(const Source& src, const std::string& msg) {
   error_ += msg;
 }
 
-void TypeDeterminer::set_referenced_from_function_if_needed(ast::Variable* var,
+void TypeDeterminer::set_referenced_from_function_if_needed(VariableInfo* var,
                                                             bool local) {
   if (current_function_ == nullptr) {
     return;
   }
-  if (var->storage_class() == ast::StorageClass::kNone ||
-      var->storage_class() == ast::StorageClass::kFunction) {
+  if (var->storage_class == ast::StorageClass::kNone ||
+      var->storage_class == ast::StorageClass::kFunction) {
     return;
   }
 
@@ -105,7 +106,18 @@ void TypeDeterminer::set_referenced_from_function_if_needed(ast::Variable* var,
 }
 
 bool TypeDeterminer::Determine() {
+  bool result = DetermineInternal();
+
+  // Even if resolving failed, create all the semantic nodes for information we
+  // did generate.
+  CreateSemanticNodes();
+
+  return result;
+}
+
+bool TypeDeterminer::DetermineInternal() {
   std::vector<type::StorageTexture*> storage_textures;
+
   for (auto& it : builder_->Types().types()) {
     if (auto* storage =
             it.second->UnwrapIfNeeded()->As<type::StorageTexture>()) {
@@ -122,7 +134,7 @@ bool TypeDeterminer::Determine() {
   }
 
   for (auto* var : builder_->AST().GlobalVariables()) {
-    variable_stack_.set_global(var->symbol(), var);
+    variable_stack_.set_global(var->symbol(), CreateVariableInfo(var));
 
     if (var->has_constructor()) {
       if (!DetermineResultType(var->constructor())) {
@@ -135,8 +147,8 @@ bool TypeDeterminer::Determine() {
     return false;
   }
 
-  // Walk over the caller to callee information and update functions with which
-  // entry points call those functions.
+  // Walk over the caller to callee information and update functions with
+  // which entry points call those functions.
   for (auto* func : builder_->AST().Functions()) {
     if (!func->IsEntryPoint()) {
       continue;
@@ -146,8 +158,6 @@ bool TypeDeterminer::Determine() {
     }
   }
 
-  CreateSemanticFunctions();
-
   return true;
 }
 
@@ -170,15 +180,13 @@ bool TypeDeterminer::DetermineFunctions(const ast::FunctionList& funcs) {
 }
 
 bool TypeDeterminer::DetermineFunction(ast::Function* func) {
-  auto* info = function_infos_.Create<FunctionInfo>(func);
-  symbol_to_function_[func->symbol()] = info;
-  function_to_info_.emplace(func, info);
-
-  current_function_ = info;
+  current_function_ = function_infos_.Create<FunctionInfo>(func);
+  symbol_to_function_[func->symbol()] = current_function_;
+  function_to_info_.emplace(func, current_function_);
 
   variable_stack_.push_scope();
   for (auto* param : func->params()) {
-    variable_stack_.set(param->symbol(), param);
+    variable_stack_.set(param->symbol(), CreateVariableInfo(param));
   }
 
   if (!DetermineStatements(func->body())) {
@@ -211,22 +219,26 @@ bool TypeDeterminer::DetermineVariableStorageClass(ast::Statement* stmt) {
   }
 
   auto* var = var_decl->variable();
+
+  auto* info = CreateVariableInfo(var);
+  variable_to_info_.emplace(var, info);
+
   // Nothing to do for const
   if (var->is_const()) {
     return true;
   }
 
-  if (var->storage_class() == ast::StorageClass::kFunction) {
+  if (info->storage_class == ast::StorageClass::kFunction) {
     return true;
   }
 
-  if (var->storage_class() != ast::StorageClass::kNone) {
+  if (info->storage_class != ast::StorageClass::kNone) {
     set_error(stmt->source(),
               "function variable has a non-function storage class");
     return false;
   }
 
-  var->set_storage_class(ast::StorageClass::kFunction);
+  info->storage_class = ast::StorageClass::kFunction;
   return true;
 }
 
@@ -291,7 +303,8 @@ bool TypeDeterminer::DetermineResultType(ast::Statement* stmt) {
     return true;
   }
   if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
-    variable_stack_.set(v->variable()->symbol(), v->variable());
+    variable_stack_.set(v->variable()->symbol(),
+                        variable_to_info_.at(v->variable()));
     return DetermineResultType(v->variable()->constructor());
   }
 
@@ -816,17 +829,17 @@ bool TypeDeterminer::DetermineConstructor(ast::ConstructorExpression* expr) {
 
 bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
   auto symbol = expr->symbol();
-  ast::Variable* var;
+  VariableInfo* var;
   if (variable_stack_.get(symbol, &var)) {
     // A constant is the type, but a variable is always a pointer so synthesize
     // the pointer around the variable type.
-    if (var->is_const()) {
-      SetType(expr, var->type());
-    } else if (var->type()->Is<type::Pointer>()) {
-      SetType(expr, var->type());
+    if (var->declaration->is_const()) {
+      SetType(expr, var->declaration->type());
+    } else if (var->declaration->type()->Is<type::Pointer>()) {
+      SetType(expr, var->declaration->type());
     } else {
-      SetType(expr, builder_->create<type::Pointer>(var->type(),
-                                                    var->storage_class()));
+      SetType(expr, builder_->create<type::Pointer>(var->declaration->type(),
+                                                    var->storage_class));
     }
 
     set_referenced_from_function_if_needed(var, true);
@@ -1146,6 +1159,13 @@ bool TypeDeterminer::DetermineUnaryOp(ast::UnaryOpExpression* expr) {
   return true;
 }
 
+TypeDeterminer::VariableInfo* TypeDeterminer::CreateVariableInfo(
+    ast::Variable* var) {
+  auto* info = variable_infos_.Create(var);
+  variable_to_info_.emplace(var, info);
+  return info;
+}
+
 bool TypeDeterminer::DetermineStorageTextureSubtype(type::StorageTexture* tex) {
   if (tex->type() != nullptr) {
     return true;
@@ -1211,22 +1231,40 @@ void TypeDeterminer::SetType(ast::Expression* expr, type::Type* type) const {
                              builder_->create<semantic::Expression>(type));
 }
 
-void TypeDeterminer::CreateSemanticFunctions() const {
+void TypeDeterminer::CreateSemanticNodes() const {
+  auto& sem = builder_->Sem();
+
+  for (auto it : variable_to_info_) {
+    auto* var = it.first;
+    auto* info = it.second;
+    sem.Add(var,
+            builder_->create<semantic::Variable>(var, info->storage_class));
+  }
+
+  auto remap_vars = [&sem](const std::vector<VariableInfo*>& in) {
+    std::vector<const semantic::Variable*> out;
+    out.reserve(in.size());
+    for (auto* info : in) {
+      out.emplace_back(sem.Get(info->declaration));
+    }
+    return out;
+  };
+
   for (auto it : function_to_info_) {
     auto* func = it.first;
     auto* info = it.second;
-    if (builder_->Sem().Get(func)) {
-      // ast::Function already has a semantic::Function node.
-      // This is likely via explicit call to DetermineXXX() in test.
-      continue;
-    }
-    builder_->Sem().Add(func, builder_->create<semantic::Function>(
-                                  info->referenced_module_vars,
-                                  info->local_referenced_module_vars,
-                                  info->ancestor_entry_points));
+    sem.Add(func, builder_->create<semantic::Function>(
+                      remap_vars(info->referenced_module_vars),
+                      remap_vars(info->local_referenced_module_vars),
+                      info->ancestor_entry_points));
   }
 }
 
+TypeDeterminer::VariableInfo::VariableInfo(ast::Variable* decl)
+    : declaration(decl), storage_class(decl->declared_storage_class()) {}
+
+TypeDeterminer::VariableInfo::~VariableInfo() = default;
+
 TypeDeterminer::FunctionInfo::FunctionInfo(ast::Function* decl)
     : declaration(decl) {}
 
diff --git a/src/type_determiner.h b/src/type_determiner.h
index c943a536b1..1402dda36a 100644
--- a/src/type_determiner.h
+++ b/src/type_determiner.h
@@ -66,23 +66,6 @@ class TypeDeterminer {
   /// @returns true if the type determiner was successful
   bool Determine();
 
-  /// Creates the semantic::Function nodes and adds them to the semantic::Info
-  /// of the ProgramBuilder.
-  void CreateSemanticFunctions() const;
-
-  /// Retrieves information for the requested import.
-  /// @param src the source of the import
-  /// @param path the import path
-  /// @param name the method name to get information on
-  /// @param params the parameters to the method call
-  /// @param id out parameter for the external call ID. Must not be a nullptr.
-  /// @returns the return type of `name` in `path` or nullptr on error.
-  type::Type* GetImportData(const Source& src,
-                            const std::string& path,
-                            const std::string& name,
-                            const ast::ExpressionList& params,
-                            uint32_t* id);
-
   /// Sets the intrinsic data information for the identifier if needed
   /// @param ident the identifier expression
   /// @returns true if an intrinsic was set
@@ -99,6 +82,7 @@ class TypeDeterminer {
         set.emplace(val);
       }
     }
+    size_t size() const { return vector.size(); }
     ConstIterator begin() const { return vector.begin(); }
     ConstIterator end() const { return vector.end(); }
     operator const std::vector<T> &() const { return vector; }
@@ -108,6 +92,16 @@ class TypeDeterminer {
     std::unordered_set<T> set;
   };
 
+  /// Structure holding semantic information about a function.
+  /// Used to build the semantic::Function nodes at the end of resolving.
+  struct VariableInfo {
+    explicit VariableInfo(ast::Variable* decl);
+    ~VariableInfo();
+
+    ast::Variable* const declaration;
+    ast::StorageClass storage_class;
+  };
+
   /// Structure holding semantic information about a function.
   /// Used to build the semantic::Function nodes at the end of resolving.
   struct FunctionInfo {
@@ -115,11 +109,16 @@ class TypeDeterminer {
     ~FunctionInfo();
 
     ast::Function* const declaration;
-    UniqueVector<ast::Variable*> referenced_module_vars;
-    UniqueVector<ast::Variable*> local_referenced_module_vars;
+    UniqueVector<VariableInfo*> referenced_module_vars;
+    UniqueVector<VariableInfo*> local_referenced_module_vars;
     UniqueVector<Symbol> ancestor_entry_points;
   };
 
+  /// Determines type information for the program, without creating final the
+  /// semantic nodes.
+  /// @returns true if the determination was successful
+  bool DetermineInternal();
+
   /// Determines type information for functions
   /// @param funcs the functions to check
   /// @returns true if the determination was successful
@@ -154,8 +153,25 @@ class TypeDeterminer {
   /// @returns false on error
   bool DetermineStorageTextureSubtype(type::StorageTexture* tex);
 
+  /// Creates the nodes and adds them to the semantic::Info mappings of the
+  /// ProgramBuilder.
+  void CreateSemanticNodes() const;
+
+  /// Retrieves information for the requested import.
+  /// @param src the source of the import
+  /// @param path the import path
+  /// @param name the method name to get information on
+  /// @param params the parameters to the method call
+  /// @param id out parameter for the external call ID. Must not be a nullptr.
+  /// @returns the return type of `name` in `path` or nullptr on error.
+  type::Type* GetImportData(const Source& src,
+                            const std::string& path,
+                            const std::string& name,
+                            const ast::ExpressionList& params,
+                            uint32_t* id);
+
   void set_error(const Source& src, const std::string& msg);
-  void set_referenced_from_function_if_needed(ast::Variable* var, bool local);
+  void set_referenced_from_function_if_needed(VariableInfo* var, bool local);
   void set_entry_points(const Symbol& fn_sym, Symbol ep_sym);
 
   bool DetermineArrayAccessor(ast::ArrayAccessorExpression* expr);
@@ -169,6 +185,8 @@ class TypeDeterminer {
   bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr);
   bool DetermineUnaryOp(ast::UnaryOpExpression* expr);
 
+  VariableInfo* CreateVariableInfo(ast::Variable*);
+
   /// @returns the resolved type of the ast::Expression `expr`
   /// @param expr the expression
   type::Type* TypeOf(ast::Expression* expr) const {
@@ -183,10 +201,12 @@ class TypeDeterminer {
 
   ProgramBuilder* builder_;
   std::string error_;
-  ScopeStack<ast::Variable*> variable_stack_;
+  ScopeStack<VariableInfo*> variable_stack_;
   std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_;
   std::unordered_map<ast::Function*, FunctionInfo*> function_to_info_;
+  std::unordered_map<ast::Variable*, VariableInfo*> variable_to_info_;
   FunctionInfo* current_function_ = nullptr;
+  BlockAllocator<VariableInfo> variable_infos_;
   BlockAllocator<FunctionInfo> function_infos_;
 
   // Map from caller functions to callee functions.
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc
index b97f009ba4..96def378b2 100644
--- a/src/type_determiner_test.cc
+++ b/src/type_determiner_test.cc
@@ -53,6 +53,7 @@
 #include "src/program_builder.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
 #include "src/type/bool_type.h"
@@ -689,11 +690,11 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables) {
 
   const auto& vars = func_sem->ReferencedModuleVariables();
   ASSERT_EQ(vars.size(), 5u);
-  EXPECT_EQ(vars[0], out_var);
-  EXPECT_EQ(vars[1], in_var);
-  EXPECT_EQ(vars[2], wg_var);
-  EXPECT_EQ(vars[3], sb_var);
-  EXPECT_EQ(vars[4], priv_var);
+  EXPECT_EQ(vars[0]->Declaration(), out_var);
+  EXPECT_EQ(vars[1]->Declaration(), in_var);
+  EXPECT_EQ(vars[2]->Declaration(), wg_var);
+  EXPECT_EQ(vars[3]->Declaration(), sb_var);
+  EXPECT_EQ(vars[4]->Declaration(), priv_var);
 }
 
 TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) {
@@ -726,11 +727,11 @@ TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) {
 
   const auto& vars = func2_sem->ReferencedModuleVariables();
   ASSERT_EQ(vars.size(), 5u);
-  EXPECT_EQ(vars[0], out_var);
-  EXPECT_EQ(vars[1], in_var);
-  EXPECT_EQ(vars[2], wg_var);
-  EXPECT_EQ(vars[3], sb_var);
-  EXPECT_EQ(vars[4], priv_var);
+  EXPECT_EQ(vars[0]->Declaration(), out_var);
+  EXPECT_EQ(vars[1]->Declaration(), in_var);
+  EXPECT_EQ(vars[2]->Declaration(), wg_var);
+  EXPECT_EQ(vars[3]->Declaration(), sb_var);
+  EXPECT_EQ(vars[4]->Declaration(), priv_var);
 }
 
 TEST_F(TypeDeterminerTest, Function_NotRegisterFunctionVariable) {
@@ -1490,7 +1491,7 @@ TEST_F(TypeDeterminerTest, StorageClass_SetsIfMissing) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  EXPECT_EQ(var->storage_class(), ast::StorageClass::kFunction);
+  EXPECT_EQ(Sem().Get(var)->StorageClass(), ast::StorageClass::kFunction);
 }
 
 TEST_F(TypeDeterminerTest, StorageClass_DoesNotSetOnConst) {
@@ -1501,7 +1502,7 @@ TEST_F(TypeDeterminerTest, StorageClass_DoesNotSetOnConst) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  EXPECT_EQ(var->storage_class(), ast::StorageClass::kNone);
+  EXPECT_EQ(Sem().Get(var)->StorageClass(), ast::StorageClass::kNone);
 }
 
 TEST_F(TypeDeterminerTest, StorageClass_NonFunctionClassError) {
diff --git a/src/validator/validator_impl.cc b/src/validator/validator_impl.cc
index 0b0b1631cd..265ad84f27 100644
--- a/src/validator/validator_impl.cc
+++ b/src/validator/validator_impl.cc
@@ -31,6 +31,7 @@
 #include "src/ast/uint_literal.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/semantic/expression.h"
+#include "src/semantic/variable.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
 #include "src/type/bool_type.h"
@@ -336,19 +337,26 @@ bool ValidatorImpl::ValidateConstructedTypes(
 bool ValidatorImpl::ValidateGlobalVariables(
     const ast::VariableList& global_vars) {
   for (auto* var : global_vars) {
+    auto* sem = program_->Sem().Get(var);
+    if (!sem) {
+      add_error(var->source(), "no semantic information for variable '" +
+                                   program_->Symbols().NameFor(var->symbol()) +
+                                   "'");
+      return false;
+    }
+
     if (variable_stack_.has(var->symbol())) {
       add_error(var->source(), "v-0011",
                 "redeclared global identifier '" +
                     program_->Symbols().NameFor(var->symbol()) + "'");
       return false;
     }
-    if (!var->is_const() && var->storage_class() == ast::StorageClass::kNone) {
+    if (!var->is_const() && sem->StorageClass() == ast::StorageClass::kNone) {
       add_error(var->source(), "v-0022",
                 "global variables must have a storage class");
       return false;
     }
-    if (var->is_const() &&
-        !(var->storage_class() == ast::StorageClass::kNone)) {
+    if (var->is_const() && !(sem->StorageClass() == ast::StorageClass::kNone)) {
       add_error(var->source(), "v-global01",
                 "global constants shouldn't have a storage class");
       return false;
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 62e52e5ce6..6e4eff6c25 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -47,6 +47,7 @@
 #include "src/program_builder.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
@@ -194,7 +195,8 @@ bool GeneratorImpl::Generate(std::ostream& out) {
 }
 
 void GeneratorImpl::register_global(ast::Variable* global) {
-  global_variables_.set(global->symbol(), global);
+  auto* sem = builder_.Sem().Get(global);
+  global_variables_.set(global->symbol(), sem);
 }
 
 std::string GeneratorImpl::generate_name(const std::string& prefix) {
@@ -1134,10 +1136,11 @@ bool GeneratorImpl::EmitExpression(std::ostream& pre,
   return false;
 }
 
-bool GeneratorImpl::global_is_in_struct(ast::Variable* var) const {
-  if (var->HasLocationDecoration() || var->HasBuiltinDecoration()) {
-    return var->storage_class() == ast::StorageClass::kInput ||
-           var->storage_class() == ast::StorageClass::kOutput;
+bool GeneratorImpl::global_is_in_struct(const semantic::Variable* var) const {
+  if (var->Declaration()->HasLocationDecoration() ||
+      var->Declaration()->HasBuiltinDecoration()) {
+    return var->StorageClass() == ast::StorageClass::kInput ||
+           var->StorageClass() == ast::StorageClass::kOutput;
   }
   return false;
 }
@@ -1146,10 +1149,10 @@ bool GeneratorImpl::EmitIdentifier(std::ostream&,
                                    std::ostream& out,
                                    ast::IdentifierExpression* expr) {
   auto* ident = expr->As<ast::IdentifierExpression>();
-  ast::Variable* var = nullptr;
+  const semantic::Variable* var = nullptr;
   if (global_variables_.get(ident->symbol(), &var)) {
     if (global_is_in_struct(var)) {
-      auto var_type = var->storage_class() == ast::StorageClass::kInput
+      auto var_type = var->StorageClass() == ast::StorageClass::kInput
                           ? VarType::kIn
                           : VarType::kOut;
       auto name = current_ep_var_name(var_type);
@@ -1230,14 +1233,14 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct(
     const semantic::Function* func) {
   for (auto data : func->ReferencedLocationVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kInput) {
+    if (var->StorageClass() == ast::StorageClass::kInput) {
       return true;
     }
   }
 
   for (auto data : func->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kInput) {
+    if (var->StorageClass() == ast::StorageClass::kInput) {
       return true;
     }
   }
@@ -1248,14 +1251,14 @@ bool GeneratorImpl::has_referenced_out_var_needing_struct(
     const semantic::Function* func) {
   for (auto data : func->ReferencedLocationVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput) {
       return true;
     }
   }
 
   for (auto data : func->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput) {
       return true;
     }
   }
@@ -1266,16 +1269,16 @@ bool GeneratorImpl::has_referenced_var_needing_struct(
     const semantic::Function* func) {
   for (auto data : func->ReferencedLocationVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput ||
-        var->storage_class() == ast::StorageClass::kInput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput ||
+        var->StorageClass() == ast::StorageClass::kInput) {
       return true;
     }
   }
 
   for (auto data : func->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput ||
-        var->storage_class() == ast::StorageClass::kInput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput ||
+        var->StorageClass() == ast::StorageClass::kInput) {
       return true;
     }
   }
@@ -1407,51 +1410,54 @@ bool GeneratorImpl::EmitEntryPointData(
 
   for (auto data : func_sem->ReferencedLocationVariables()) {
     auto* var = data.first;
+    auto* decl = var->Declaration();
     auto* deco = data.second;
 
-    if (var->storage_class() == ast::StorageClass::kInput) {
-      in_variables.push_back({var, deco});
-    } else if (var->storage_class() == ast::StorageClass::kOutput) {
-      outvariables.push_back({var, deco});
+    if (var->StorageClass() == ast::StorageClass::kInput) {
+      in_variables.push_back({decl, deco});
+    } else if (var->StorageClass() == ast::StorageClass::kOutput) {
+      outvariables.push_back({decl, deco});
     }
   }
 
   for (auto data : func_sem->ReferencedBuiltinVariables()) {
     auto* var = data.first;
+    auto* decl = var->Declaration();
     auto* deco = data.second;
 
-    if (var->storage_class() == ast::StorageClass::kInput) {
-      in_variables.push_back({var, deco});
-    } else if (var->storage_class() == ast::StorageClass::kOutput) {
-      outvariables.push_back({var, deco});
+    if (var->StorageClass() == ast::StorageClass::kInput) {
+      in_variables.push_back({decl, deco});
+    } else if (var->StorageClass() == ast::StorageClass::kOutput) {
+      outvariables.push_back({decl, deco});
     }
   }
 
   bool emitted_uniform = false;
   for (auto data : func_sem->ReferencedUniformVariables()) {
     auto* var = data.first;
+    auto* decl = var->Declaration();
     // TODO(dsinclair): We're using the binding to make up the buffer number but
     // we should instead be using a provided mapping that uses both buffer and
     // set. https://bugs.chromium.org/p/tint/issues/detail?id=104
     auto* binding = data.second.binding;
     if (binding == nullptr) {
       error_ = "unable to find binding information for uniform: " +
-               builder_.Symbols().NameFor(var->symbol());
+               builder_.Symbols().NameFor(decl->symbol());
       return false;
     }
     // auto* set = data.second.set;
 
     // If the global has already been emitted we skip it, it's been emitted by
     // a previous entry point.
-    if (emitted_globals.count(var->symbol()) != 0) {
+    if (emitted_globals.count(decl->symbol()) != 0) {
       continue;
     }
-    emitted_globals.insert(var->symbol());
+    emitted_globals.insert(decl->symbol());
 
-    auto* type = var->type()->UnwrapIfNeeded();
+    auto* type = decl->type()->UnwrapIfNeeded();
     if (auto* strct = type->As<type::Struct>()) {
       out << "ConstantBuffer<" << builder_.Symbols().NameFor(strct->symbol())
-          << "> " << builder_.Symbols().NameFor(var->symbol())
+          << "> " << builder_.Symbols().NameFor(decl->symbol())
           << " : register(b" << binding->value() << ");" << std::endl;
     } else {
       // TODO(dsinclair): There is outstanding spec work to require all uniform
@@ -1460,7 +1466,7 @@ bool GeneratorImpl::EmitEntryPointData(
       // is not a block.
       // Relevant: https://github.com/gpuweb/gpuweb/issues/1004
       //           https://github.com/gpuweb/gpuweb/issues/1008
-      auto name = "cbuffer_" + builder_.Symbols().NameFor(var->symbol());
+      auto name = "cbuffer_" + builder_.Symbols().NameFor(decl->symbol());
       out << "cbuffer " << name << " : register(b" << binding->value() << ") {"
           << std::endl;
 
@@ -1469,7 +1475,7 @@ bool GeneratorImpl::EmitEntryPointData(
       if (!EmitType(out, type, "")) {
         return false;
       }
-      out << " " << builder_.Symbols().NameFor(var->symbol()) << ";"
+      out << " " << builder_.Symbols().NameFor(decl->symbol()) << ";"
           << std::endl;
       decrement_indent();
       out << "};" << std::endl;
@@ -1484,16 +1490,17 @@ bool GeneratorImpl::EmitEntryPointData(
   bool emitted_storagebuffer = false;
   for (auto data : func_sem->ReferencedStoragebufferVariables()) {
     auto* var = data.first;
+    auto* decl = var->Declaration();
     auto* binding = data.second.binding;
 
     // If the global has already been emitted we skip it, it's been emitted by
     // a previous entry point.
-    if (emitted_globals.count(var->symbol()) != 0) {
+    if (emitted_globals.count(decl->symbol()) != 0) {
       continue;
     }
-    emitted_globals.insert(var->symbol());
+    emitted_globals.insert(decl->symbol());
 
-    auto* ac = var->type()->As<type::AccessControl>();
+    auto* ac = decl->type()->As<type::AccessControl>();
     if (ac == nullptr) {
       error_ = "access control type required for storage buffer";
       return false;
@@ -1502,7 +1509,7 @@ bool GeneratorImpl::EmitEntryPointData(
     if (ac->IsReadWrite()) {
       out << "RW";
     }
-    out << "ByteAddressBuffer " << builder_.Symbols().NameFor(var->symbol())
+    out << "ByteAddressBuffer " << builder_.Symbols().NameFor(decl->symbol())
         << " : register(u" << binding->value() << ");" << std::endl;
     emitted_storagebuffer = true;
   }
@@ -2069,11 +2076,11 @@ bool GeneratorImpl::is_storage_buffer_access(
 
   // Check if this is a storage buffer variable
   if (auto* ident = expr->structure()->As<ast::IdentifierExpression>()) {
-    ast::Variable* var = nullptr;
+    const semantic::Variable* var = nullptr;
     if (!global_variables_.get(ident->symbol(), &var)) {
       return false;
     }
-    return var->storage_class() == ast::StorageClass::kStorage;
+    return var->StorageClass() == ast::StorageClass::kStorage;
   } else if (auto* member = structure->As<ast::MemberAccessorExpression>()) {
     return is_storage_buffer_access(member);
   } else if (auto* array = structure->As<ast::ArrayAccessorExpression>()) {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 24089ae64f..0c175fc86c 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -338,7 +338,7 @@ class GeneratorImpl {
   /// Checks if the global variable is in an input or output struct
   /// @param var the variable to check
   /// @returns true if the global is in an input or output struct
-  bool global_is_in_struct(ast::Variable* var) const;
+  bool global_is_in_struct(const semantic::Variable* var) const;
   /// Creates a text string representing the index into a storage buffer
   /// @param pre the pre stream
   /// @param expr the expression to use as the index
@@ -400,7 +400,7 @@ class GeneratorImpl {
   Symbol current_ep_sym_;
   bool generating_entry_point_ = false;
   uint32_t loop_emission_counter_ = 0;
-  ScopeStack<ast::Variable*> global_variables_;
+  ScopeStack<const semantic::Variable*> global_variables_;
   std::unordered_map<Symbol, EntryPointData> ep_sym_to_in_data_;
   std::unordered_map<Symbol, EntryPointData> ep_sym_to_out_data_;
 
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 0f9d249679..be67f4d96c 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -52,6 +52,7 @@
 #include "src/program.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
@@ -120,7 +121,8 @@ bool GeneratorImpl::Generate() {
   out_ << "#include <metal_stdlib>" << std::endl << std::endl;
 
   for (auto* global : program_->AST().GlobalVariables()) {
-    global_variables_.set(global->symbol(), global);
+    auto* sem = program_->Sem().Get(global);
+    global_variables_.set(global->symbol(), sem);
   }
 
   for (auto* const ty : program_->AST().ConstructedTypes()) {
@@ -508,14 +510,14 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
   auto* func_sem = program_->Sem().Get(func);
   for (const auto& data : func_sem->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() != ast::StorageClass::kInput) {
+    if (var->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
     if (!first) {
       out_ << ", ";
     }
     first = false;
-    out_ << program_->Symbols().NameFor(var->symbol());
+    out_ << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   for (const auto& data : func_sem->ReferencedUniformVariables()) {
@@ -524,7 +526,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << ", ";
     }
     first = false;
-    out_ << program_->Symbols().NameFor(var->symbol());
+    out_ << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   for (const auto& data : func_sem->ReferencedStoragebufferVariables()) {
@@ -533,7 +535,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << ", ";
     }
     first = false;
-    out_ << program_->Symbols().NameFor(var->symbol());
+    out_ << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   const auto& params = expr->params();
@@ -1033,10 +1035,10 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
     auto* var = data.first;
     auto* deco = data.second;
 
-    if (var->storage_class() == ast::StorageClass::kInput) {
-      in_locations.push_back({var, deco->value()});
-    } else if (var->storage_class() == ast::StorageClass::kOutput) {
-      out_variables.push_back({var, deco});
+    if (var->StorageClass() == ast::StorageClass::kInput) {
+      in_locations.push_back({var->Declaration(), deco->value()});
+    } else if (var->StorageClass() == ast::StorageClass::kOutput) {
+      out_variables.push_back({var->Declaration(), deco});
     }
   }
 
@@ -1044,8 +1046,8 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
     auto* var = data.first;
     auto* deco = data.second;
 
-    if (var->storage_class() == ast::StorageClass::kOutput) {
-      out_variables.push_back({var, deco});
+    if (var->StorageClass() == ast::StorageClass::kOutput) {
+      out_variables.push_back({var->Declaration(), deco});
     }
   }
 
@@ -1191,7 +1193,7 @@ bool GeneratorImpl::has_referenced_in_var_needing_struct(ast::Function* func) {
   auto* func_sem = program_->Sem().Get(func);
   for (auto data : func_sem->ReferencedLocationVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kInput) {
+    if (var->StorageClass() == ast::StorageClass::kInput) {
       return true;
     }
   }
@@ -1203,14 +1205,14 @@ bool GeneratorImpl::has_referenced_out_var_needing_struct(ast::Function* func) {
 
   for (auto data : func_sem->ReferencedLocationVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput) {
       return true;
     }
   }
 
   for (auto data : func_sem->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() == ast::StorageClass::kOutput) {
+    if (var->StorageClass() == ast::StorageClass::kOutput) {
       return true;
     }
   }
@@ -1308,7 +1310,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
 
   for (const auto& data : func_sem->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() != ast::StorageClass::kInput) {
+    if (var->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
     if (!first) {
@@ -1317,10 +1319,10 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     first = false;
 
     out_ << "thread ";
-    if (!EmitType(var->type(), "")) {
+    if (!EmitType(var->Declaration()->type(), "")) {
       return false;
     }
-    out_ << "& " << program_->Symbols().NameFor(var->symbol());
+    out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   for (const auto& data : func_sem->ReferencedUniformVariables()) {
@@ -1332,10 +1334,10 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
 
     out_ << "constant ";
     // TODO(dsinclair): Can arrays be uniform? If so, fix this ...
-    if (!EmitType(var->type(), "")) {
+    if (!EmitType(var->Declaration()->type(), "")) {
       return false;
     }
-    out_ << "& " << program_->Symbols().NameFor(var->symbol());
+    out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   for (const auto& data : func_sem->ReferencedStoragebufferVariables()) {
@@ -1345,7 +1347,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     }
     first = false;
 
-    auto* ac = var->type()->As<type::AccessControl>();
+    auto* ac = var->Declaration()->type()->As<type::AccessControl>();
     if (ac == nullptr) {
       error_ = "invalid type for storage buffer, expected access control";
       return false;
@@ -1358,7 +1360,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     if (!EmitType(ac->type(), "")) {
       return false;
     }
-    out_ << "& " << program_->Symbols().NameFor(var->symbol());
+    out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol());
   }
 
   for (auto* v : func->params()) {
@@ -1447,7 +1449,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
 
   for (auto data : func_sem->ReferencedBuiltinVariables()) {
     auto* var = data.first;
-    if (var->storage_class() != ast::StorageClass::kInput) {
+    if (var->StorageClass() != ast::StorageClass::kInput) {
       continue;
     }
 
@@ -1458,7 +1460,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
 
     auto* builtin = data.second;
 
-    if (!EmitType(var->type(), "")) {
+    if (!EmitType(var->Declaration()->type(), "")) {
       return false;
     }
 
@@ -1467,8 +1469,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
       error_ = "unknown builtin";
       return false;
     }
-    out_ << " " << program_->Symbols().NameFor(var->symbol()) << " [[" << attr
-         << "]]";
+    out_ << " " << program_->Symbols().NameFor(var->Declaration()->symbol())
+         << " [[" << attr << "]]";
   }
 
   for (auto data : func_sem->ReferencedUniformVariables()) {
@@ -1484,7 +1486,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     auto* binding = data.second.binding;
     if (binding == nullptr) {
       error_ = "unable to find binding information for uniform: " +
-               program_->Symbols().NameFor(var->symbol());
+               program_->Symbols().NameFor(var->Declaration()->symbol());
       return false;
     }
     // auto* set = data.second.set;
@@ -1492,11 +1494,11 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     out_ << "constant ";
     // TODO(dsinclair): Can you have a uniform array? If so, this needs to be
     // updated to handle arrays property.
-    if (!EmitType(var->type(), "")) {
+    if (!EmitType(var->Declaration()->type(), "")) {
       return false;
     }
-    out_ << "& " << program_->Symbols().NameFor(var->symbol()) << " [[buffer("
-         << binding->value() << ")]]";
+    out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol())
+         << " [[buffer(" << binding->value() << ")]]";
   }
 
   for (auto data : func_sem->ReferencedStoragebufferVariables()) {
@@ -1512,7 +1514,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     auto* binding = data.second.binding;
     // auto* set = data.second.set;
 
-    auto* ac = var->type()->As<type::AccessControl>();
+    auto* ac = var->Declaration()->type()->As<type::AccessControl>();
     if (ac == nullptr) {
       error_ = "invalid type for storage buffer, expected access control";
       return false;
@@ -1525,8 +1527,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     if (!EmitType(ac->type(), "")) {
       return false;
     }
-    out_ << "& " << program_->Symbols().NameFor(var->symbol()) << " [[buffer("
-         << binding->value() << ")]]";
+    out_ << "& " << program_->Symbols().NameFor(var->Declaration()->symbol())
+         << " [[buffer(" << binding->value() << ")]]";
   }
 
   out_ << ") {" << std::endl;
@@ -1563,23 +1565,23 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
   return true;
 }
 
-bool GeneratorImpl::global_is_in_struct(ast::Variable* var) const {
+bool GeneratorImpl::global_is_in_struct(const semantic::Variable* var) const {
   bool in_or_out_struct_has_location =
-      var != nullptr && var->HasLocationDecoration() &&
-      (var->storage_class() == ast::StorageClass::kInput ||
-       var->storage_class() == ast::StorageClass::kOutput);
+      var != nullptr && var->Declaration()->HasLocationDecoration() &&
+      (var->StorageClass() == ast::StorageClass::kInput ||
+       var->StorageClass() == ast::StorageClass::kOutput);
   bool in_struct_has_builtin =
-      var != nullptr && var->HasBuiltinDecoration() &&
-      var->storage_class() == ast::StorageClass::kOutput;
+      var != nullptr && var->Declaration()->HasBuiltinDecoration() &&
+      var->StorageClass() == ast::StorageClass::kOutput;
   return in_or_out_struct_has_location || in_struct_has_builtin;
 }
 
 bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) {
   auto* ident = expr->As<ast::IdentifierExpression>();
-  ast::Variable* var = nullptr;
+  const semantic::Variable* var = nullptr;
   if (global_variables_.get(ident->symbol(), &var)) {
     if (global_is_in_struct(var)) {
-      auto var_type = var->storage_class() == ast::StorageClass::kInput
+      auto var_type = var->StorageClass() == ast::StorageClass::kInput
                           ? VarType::kIn
                           : VarType::kOut;
       auto name = current_ep_var_name(var_type);
@@ -1623,7 +1625,7 @@ bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
     // will be turned into assignments.
     for (auto* s : *(stmt->body())) {
       if (auto* decl = s->As<ast::VariableDeclStatement>()) {
-        if (!EmitVariable(decl->variable(), true)) {
+        if (!EmitVariable(program_->Sem().Get(decl->variable()), true)) {
           return false;
         }
       }
@@ -1839,7 +1841,8 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
     return EmitSwitch(s);
   }
   if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
-    return EmitVariable(v->variable(), false);
+    auto* var = program_->Sem().Get(v->variable());
+    return EmitVariable(var, false);
   }
 
   error_ = "unknown statement type: " + program_->str(stmt);
@@ -2078,35 +2081,38 @@ bool GeneratorImpl::EmitUnaryOp(ast::UnaryOpExpression* expr) {
   return true;
 }
 
-bool GeneratorImpl::EmitVariable(ast::Variable* var, bool skip_constructor) {
+bool GeneratorImpl::EmitVariable(const semantic::Variable* var,
+                                 bool skip_constructor) {
   make_indent();
 
+  auto* decl = var->Declaration();
+
   // TODO(dsinclair): Handle variable decorations
-  if (!var->decorations().empty()) {
+  if (!decl->decorations().empty()) {
     error_ = "Variable decorations are not handled yet";
     return false;
   }
-  if (var->is_const()) {
+  if (decl->is_const()) {
     out_ << "const ";
   }
-  if (!EmitType(var->type(), program_->Symbols().NameFor(var->symbol()))) {
+  if (!EmitType(decl->type(), program_->Symbols().NameFor(decl->symbol()))) {
     return false;
   }
-  if (!var->type()->Is<type::Array>()) {
-    out_ << " " << program_->Symbols().NameFor(var->symbol());
+  if (!decl->type()->Is<type::Array>()) {
+    out_ << " " << program_->Symbols().NameFor(decl->symbol());
   }
 
   if (!skip_constructor) {
     out_ << " = ";
-    if (var->constructor() != nullptr) {
-      if (!EmitExpression(var->constructor())) {
+    if (decl->constructor() != nullptr) {
+      if (!EmitExpression(decl->constructor())) {
         return false;
       }
-    } else if (var->storage_class() == ast::StorageClass::kPrivate ||
-               var->storage_class() == ast::StorageClass::kFunction ||
-               var->storage_class() == ast::StorageClass::kNone ||
-               var->storage_class() == ast::StorageClass::kOutput) {
-      if (!EmitZeroValue(var->type())) {
+    } else if (var->StorageClass() == ast::StorageClass::kPrivate ||
+               var->StorageClass() == ast::StorageClass::kFunction ||
+               var->StorageClass() == ast::StorageClass::kNone ||
+               var->StorageClass() == ast::StorageClass::kOutput) {
+      if (!EmitZeroValue(decl->type())) {
         return false;
       }
     }
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index d67b54dcda..0aafa6261e 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -221,7 +221,7 @@ class GeneratorImpl : public TextGenerator {
   /// @param var the variable to generate
   /// @param skip_constructor set true if the constructor should be skipped
   /// @returns true if the variable was emitted
-  bool EmitVariable(ast::Variable* var, bool skip_constructor);
+  bool EmitVariable(const semantic::Variable* var, bool skip_constructor);
   /// Handles generating a program scope constant variable
   /// @param var the variable to emit
   /// @returns true if the variable was emitted
@@ -256,7 +256,7 @@ class GeneratorImpl : public TextGenerator {
   /// Checks if the global variable is in an input or output struct
   /// @param var the variable to check
   /// @returns true if the global is in an input or output struct
-  bool global_is_in_struct(ast::Variable* var) const;
+  bool global_is_in_struct(const semantic::Variable* var) const;
 
   /// Converts a builtin to an attribute name
   /// @param builtin the builtin to convert
@@ -283,7 +283,7 @@ class GeneratorImpl : public TextGenerator {
   }
 
   Namer namer_;
-  ScopeStack<ast::Variable*> global_variables_;
+  ScopeStack<const semantic::Variable*> global_variables_;
   Symbol current_ep_sym_;
   bool generating_entry_point_ = false;
   const Program* program_ = nullptr;
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 645c6b4e1b..79c71e864b 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -61,6 +61,7 @@
 #include "src/program.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
@@ -462,15 +463,15 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
   for (const auto* var : func_sem->ReferencedModuleVariables()) {
     // For SPIR-V 1.3 we only output Input/output variables. If we update to
     // SPIR-V 1.4 or later this should be all variables.
-    if (var->storage_class() != ast::StorageClass::kInput &&
-        var->storage_class() != ast::StorageClass::kOutput) {
+    if (var->StorageClass() != ast::StorageClass::kInput &&
+        var->StorageClass() != ast::StorageClass::kOutput) {
       continue;
     }
 
     uint32_t var_id;
-    if (!scope_stack_.get(var->symbol(), &var_id)) {
+    if (!scope_stack_.get(var->Declaration()->symbol(), &var_id)) {
       error_ = "unable to find ID for global variable: " +
-               builder_.Symbols().NameFor(var->symbol());
+               builder_.Symbols().NameFor(var->Declaration()->symbol());
       return false;
     }
 
@@ -700,6 +701,8 @@ bool Builder::GenerateStore(uint32_t to, uint32_t from) {
 }
 
 bool Builder::GenerateGlobalVariable(ast::Variable* var) {
+  auto* sem = builder_.Sem().Get(var);
+
   uint32_t init_id = 0;
   if (var->has_constructor()) {
     if (!var->constructor()->Is<ast::ConstructorExpression>()) {
@@ -731,9 +734,9 @@ bool Builder::GenerateGlobalVariable(ast::Variable* var) {
   auto result = result_op();
   auto var_id = result.to_i();
 
-  auto sc = var->storage_class() == ast::StorageClass::kNone
+  auto sc = sem->StorageClass() == ast::StorageClass::kNone
                 ? ast::StorageClass::kPrivate
-                : var->storage_class();
+                : sem->StorageClass();
 
   type::Pointer pt(var->type(), sc);
   auto type_id = GenerateTypeIfNeeded(&pt);
@@ -799,9 +802,9 @@ bool Builder::GenerateGlobalVariable(ast::Variable* var) {
         return 0;
       }
       ops.push_back(Operand::Int(init_id));
-    } else if (var->storage_class() == ast::StorageClass::kPrivate ||
-               var->storage_class() == ast::StorageClass::kNone ||
-               var->storage_class() == ast::StorageClass::kOutput) {
+    } else if (sem->StorageClass() == ast::StorageClass::kPrivate ||
+               sem->StorageClass() == ast::StorageClass::kNone ||
+               sem->StorageClass() == ast::StorageClass::kOutput) {
       ast::NullLiteral nl(Source{}, type);
       init_id = GenerateLiteralIfNeeded(var, &nl);
       if (init_id == 0) {
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 7410675778..3025ffddc9 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -59,6 +59,7 @@
 #include "src/ast/workgroup_decoration.h"
 #include "src/program.h"
 #include "src/semantic/function.h"
+#include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
@@ -148,7 +149,7 @@ bool GeneratorImpl::GenerateEntryPoint(ast::PipelineStage stage,
 
   bool found_func_variable = false;
   for (auto* var : program_->Sem().Get(func)->ReferencedModuleVariables()) {
-    if (!EmitVariable(var)) {
+    if (!EmitVariable(var->Declaration())) {
       return false;
     }
     found_func_variable = true;
@@ -583,9 +584,11 @@ bool GeneratorImpl::EmitStructType(const type::Struct* str) {
 }
 
 bool GeneratorImpl::EmitVariable(ast::Variable* var) {
+  auto* sem = program_->Sem().Get(var);
+
   make_indent();
 
-  if (!var->decorations().empty() && !EmitVariableDecorations(var)) {
+  if (!var->decorations().empty() && !EmitVariableDecorations(sem)) {
     return false;
   }
 
@@ -593,9 +596,9 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) {
     out_ << "const";
   } else {
     out_ << "var";
-    if (var->storage_class() != ast::StorageClass::kNone &&
-        var->storage_class() != ast::StorageClass::kFunction) {
-      out_ << "<" << var->storage_class() << ">";
+    if (sem->StorageClass() != ast::StorageClass::kNone &&
+        sem->StorageClass() != ast::StorageClass::kFunction) {
+      out_ << "<" << sem->StorageClass() << ">";
     }
   }
 
@@ -615,10 +618,12 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) {
   return true;
 }
 
-bool GeneratorImpl::EmitVariableDecorations(ast::Variable* var) {
+bool GeneratorImpl::EmitVariableDecorations(const semantic::Variable* var) {
+  auto* decl = var->Declaration();
+
   out_ << "[[";
   bool first = true;
-  for (auto* deco : var->decorations()) {
+  for (auto* deco : decl->decorations()) {
     if (!first) {
       out_ << ", ";
     }
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h
index fa4603b3f2..34ad398261 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -202,7 +202,7 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating variable decorations
   /// @param var the decorated variable
   /// @returns true if the variable decoration was emitted
-  bool EmitVariableDecorations(ast::Variable* var);
+  bool EmitVariableDecorations(const semantic::Variable* var);
 
  private:
   Program const* const program_;