diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 5594bf8..980b16d 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -377,6 +377,8 @@
     "resolver/resolver.h",
     "resolver/resolver_constants.cc",
     "resolver/resolver_validation.cc",
+    "resolver/sem_helper.cc",
+    "resolver/sem_helper.h",
     "scope_stack.h",
     "sem/array.h",
     "sem/atomic_type.h",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 527eeaf..ae3f224 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -258,6 +258,8 @@
   resolver/resolver_constants.cc
   resolver/resolver_validation.cc
   resolver/resolver.h
+  resolver/sem_helper.cc
+  resolver/sem_helper.h
   scope_stack.h
   sem/array.cc
   sem/array.h
diff --git a/src/tint/bench/benchmark.h b/src/tint/bench/benchmark.h
index 96a935d..22605b2 100644
--- a/src/tint/bench/benchmark.h
+++ b/src/tint/bench/benchmark.h
@@ -17,6 +17,7 @@
 
 #include <memory>
 #include <string>
+// TODO(https://crbug.com/dawn/1379) Update cpplint and remove NOLINT
 #include <variant>  // NOLINT: Found C system header after C++ system header.
 
 #include "benchmark/benchmark.h"
diff --git a/src/tint/reader/spirv/parser_impl_module_var_test.cc b/src/tint/reader/spirv/parser_impl_module_var_test.cc
index d50959d..c8f6935 100644
--- a/src/tint/reader/spirv/parser_impl_module_var_test.cc
+++ b/src/tint/reader/spirv/parser_impl_module_var_test.cc
@@ -1272,7 +1272,6 @@
       << module_str;
 }
 
-
 TEST_F(SpvModuleScopeVarParserTest, BindingDecoration_Valid) {
   auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
      OpDecorate %1 DescriptorSet 0 ; WGSL validation requires this already
diff --git a/src/tint/reader/wgsl/token.h b/src/tint/reader/wgsl/token.h
index 7b5b675..a3715ee 100644
--- a/src/tint/reader/wgsl/token.h
+++ b/src/tint/reader/wgsl/token.h
@@ -17,7 +17,8 @@
 
 #include <string>
 #include <string_view>
-#include <variant>  // NOLINT: cpplint doesn't recognise this
+// TODO(https://crbug.com/dawn/1379) Update cpplint and remove NOLINT
+#include <variant>  // NOLINT(build/include_order))
 
 #include "src/tint/source.h"
 
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 9d125e7..4f7c08d 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -84,7 +84,8 @@
 Resolver::Resolver(ProgramBuilder* builder)
     : builder_(builder),
       diagnostics_(builder->Diagnostics()),
-      builtin_table_(BuiltinTable::Create(*builder)) {}
+      builtin_table_(BuiltinTable::Create(*builder)),
+      sem_(builder) {}
 
 Resolver::~Resolver() = default;
 
@@ -505,7 +506,7 @@
       next_constant_id = constant_id + 1;
     }
 
-    auto* sem = Sem<sem::GlobalVariable>(var);
+    auto* sem = sem_.Get<sem::GlobalVariable>(var);
     const_cast<sem::GlobalVariable*>(sem)->SetConstantId(constant_id);
   }
 }
@@ -513,9 +514,11 @@
 void Resolver::SetShadows() {
   for (auto it : dependencies_.shadows) {
     Switch(
-        Sem(it.first),  //
-        [&](sem::LocalVariable* local) { local->SetShadows(Sem(it.second)); },
-        [&](sem::Parameter* param) { param->SetShadows(Sem(it.second)); });
+        sem_.Get(it.first),  //
+        [&](sem::LocalVariable* local) {
+          local->SetShadows(sem_.Get(it.second));
+        },
+        [&](sem::Parameter* param) { param->SetShadows(sem_.Get(it.second)); });
   }
 }
 
@@ -554,7 +557,7 @@
 
   // TODO(bclayton): Call this at the end of resolve on all uniform and storage
   // referenced structs
-  if (!ValidateStorageClassLayout(sem)) {
+  if (!ValidateStorageClassLayout(sem, valid_type_storage_layouts_)) {
     return nullptr;
   }
 
@@ -699,7 +702,10 @@
     return nullptr;
   }
 
-  if (!ValidateFunction(func)) {
+  auto stage = current_function_
+                   ? current_function_->Declaration()->PipelineStage()
+                   : ast::PipelineStage::kNone;
+  if (!ValidateFunction(func, stage)) {
     return nullptr;
   }
 
@@ -752,7 +758,7 @@
         "workgroup_size arguments must be of the same type, either i32 "
         "or u32";
 
-    auto* ty = TypeOf(expr);
+    auto* ty = sem_.TypeOf(expr);
     bool is_i32 = ty->UnwrapRef()->Is<sem::I32>();
     bool is_u32 = ty->UnwrapRef()->Is<sem::U32>();
     if (!is_i32 && !is_u32) {
@@ -769,7 +775,7 @@
 
     sem::Constant value;
 
-    if (auto* user = Sem(expr)->As<sem::VariableUser>()) {
+    if (auto* user = sem_.Get(expr)->As<sem::VariableUser>()) {
       // We have an variable of a module-scope constant.
       auto* decl = user->Variable()->Declaration();
       if (!decl->is_const) {
@@ -782,14 +788,14 @@
       }
 
       if (decl->constructor) {
-        value = Sem(decl->constructor)->ConstantValue();
+        value = sem_.Get(decl->constructor)->ConstantValue();
       } else {
         // No constructor means this value must be overriden by the user.
         ws[i].value = 0;
         continue;
       }
     } else if (expr->Is<ast::LiteralExpression>()) {
-      value = Sem(expr)->ConstantValue();
+      value = sem_.Get(expr)->ConstantValue();
     } else {
       AddError(
           "workgroup_size argument must be either a literal or a "
@@ -1165,8 +1171,8 @@
 
 sem::Expression* Resolver::IndexAccessor(
     const ast::IndexAccessorExpression* expr) {
-  auto* idx = Sem(expr->index);
-  auto* obj = Sem(expr->object);
+  auto* idx = sem_.Get(expr->index);
+  auto* obj = sem_.Get(expr->object);
   auto* obj_raw_ty = obj->Type();
   auto* obj_ty = obj_raw_ty->UnwrapRef();
   auto* ty = Switch(
@@ -1177,7 +1183,7 @@
         return builder_->create<sem::Vector>(mat->type(), mat->rows());
       },
       [&](Default) {
-        AddError("cannot index type '" + TypeNameOf(obj_ty) + "'",
+        AddError("cannot index type '" + sem_.TypeNameOf(obj_ty) + "'",
                  expr->source);
         return nullptr;
       });
@@ -1188,7 +1194,7 @@
   auto* idx_ty = idx->Type()->UnwrapRef();
   if (!idx_ty->IsAnyOf<sem::I32, sem::U32>()) {
     AddError("index must be of type 'i32' or 'u32', found: '" +
-                 TypeNameOf(idx_ty) + "'",
+                 sem_.TypeNameOf(idx_ty) + "'",
              idx->Declaration()->source);
     return nullptr;
   }
@@ -1208,7 +1214,7 @@
 }
 
 sem::Expression* Resolver::Bitcast(const ast::BitcastExpression* expr) {
-  auto* inner = Sem(expr->expr);
+  auto* inner = sem_.Get(expr->expr);
   auto* ty = Type(expr->type);
   if (!ty) {
     return nullptr;
@@ -1237,7 +1243,7 @@
   const sem::Type* arg_el_ty = nullptr;
 
   for (size_t i = 0; i < expr->args.size(); i++) {
-    auto* arg = Sem(expr->args[i]);
+    auto* arg = sem_.Get(expr->args[i]);
     if (!arg) {
       return nullptr;
     }
@@ -1634,7 +1640,7 @@
 }
 
 sem::Expression* Resolver::Literal(const ast::LiteralExpression* literal) {
-  auto* ty = TypeOf(literal);
+  auto* ty = sem_.TypeOf(literal);
   if (!ty) {
     return nullptr;
   }
@@ -1722,14 +1728,14 @@
 
 sem::Expression* Resolver::MemberAccessor(
     const ast::MemberAccessorExpression* expr) {
-  auto* structure = TypeOf(expr->structure);
+  auto* structure = sem_.TypeOf(expr->structure);
   auto* storage_ty = structure->UnwrapRef();
 
   const sem::Type* ret = nullptr;
   std::vector<uint32_t> swizzle;
 
   // Structure may be a side-effecting expression (e.g. function call).
-  auto* sem_structure = Sem(expr->structure);
+  auto* sem_structure = sem_.Get(expr->structure);
   bool has_side_effects = sem_structure && sem_structure->HasSideEffects();
 
   if (auto* str = storage_ty->As<sem::Struct>()) {
@@ -1837,14 +1843,14 @@
 
   AddError(
       "invalid member accessor expression. Expected vector or struct, got '" +
-          TypeNameOf(storage_ty) + "'",
+          sem_.TypeNameOf(storage_ty) + "'",
       expr->structure->source);
   return nullptr;
 }
 
 sem::Expression* Resolver::Binary(const ast::BinaryExpression* expr) {
-  auto* lhs = Sem(expr->lhs);
-  auto* rhs = Sem(expr->rhs);
+  auto* lhs = sem_.Get(expr->lhs);
+  auto* rhs = sem_.Get(expr->rhs);
   auto* lhs_ty = lhs->Type()->UnwrapRef();
   auto* rhs_ty = rhs->Type()->UnwrapRef();
 
@@ -1852,8 +1858,8 @@
   if (!ty) {
     AddError(
         "Binary expression operand types are invalid for this operation: " +
-            TypeNameOf(lhs_ty) + " " + FriendlyName(expr->op) + " " +
-            TypeNameOf(rhs_ty),
+            sem_.TypeNameOf(lhs_ty) + " " + FriendlyName(expr->op) + " " +
+            sem_.TypeNameOf(rhs_ty),
         expr->source);
     return nullptr;
   }
@@ -2033,7 +2039,7 @@
 }
 
 sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
-  auto* expr = Sem(unary->expr);
+  auto* expr = sem_.Get(unary->expr);
   auto* expr_ty = expr->Type();
   if (!expr_ty) {
     return nullptr;
@@ -2046,9 +2052,9 @@
       // Result type matches the deref'd inner type.
       ty = expr_ty->UnwrapRef();
       if (!ty->Is<sem::Bool>() && !ty->is_bool_vector()) {
-        AddError(
-            "cannot logical negate expression of type '" + TypeNameOf(expr_ty),
-            unary->expr->source);
+        AddError("cannot logical negate expression of type '" +
+                     sem_.TypeNameOf(expr_ty),
+                 unary->expr->source);
         return nullptr;
       }
       break;
@@ -2058,7 +2064,7 @@
       ty = expr_ty->UnwrapRef();
       if (!ty->is_integer_scalar_or_vector()) {
         AddError("cannot bitwise complement expression of type '" +
-                     TypeNameOf(expr_ty),
+                     sem_.TypeNameOf(expr_ty),
                  unary->expr->source);
         return nullptr;
       }
@@ -2069,8 +2075,9 @@
       ty = expr_ty->UnwrapRef();
       if (!(ty->IsAnyOf<sem::F32, sem::I32>() ||
             ty->is_signed_integer_vector() || ty->is_float_vector())) {
-        AddError("cannot negate expression of type '" + TypeNameOf(expr_ty),
-                 unary->expr->source);
+        AddError(
+            "cannot negate expression of type '" + sem_.TypeNameOf(expr_ty),
+            unary->expr->source);
         return nullptr;
       }
       break;
@@ -2086,9 +2093,10 @@
 
         auto* array = unary->expr->As<ast::IndexAccessorExpression>();
         auto* member = unary->expr->As<ast::MemberAccessorExpression>();
-        if ((array && TypeOf(array->object)->UnwrapRef()->Is<sem::Vector>()) ||
+        if ((array &&
+             sem_.TypeOf(array->object)->UnwrapRef()->Is<sem::Vector>()) ||
             (member &&
-             TypeOf(member->structure)->UnwrapRef()->Is<sem::Vector>())) {
+             sem_.TypeOf(member->structure)->UnwrapRef()->Is<sem::Vector>())) {
           AddError("cannot take the address of a vector component",
                    unary->expr->source);
           return nullptr;
@@ -2108,7 +2116,7 @@
             ptr->StoreType(), ptr->StorageClass(), ptr->Access());
       } else {
         AddError("cannot dereference expression of type '" +
-                     TypeNameOf(expr_ty) + "'",
+                     sem_.TypeNameOf(expr_ty) + "'",
                  unary->expr->source);
         return nullptr;
       }
@@ -2140,41 +2148,6 @@
   return result;
 }
 
-sem::Type* Resolver::TypeOf(const ast::Expression* expr) {
-  auto* sem = Sem(expr);
-  return sem ? const_cast<sem::Type*>(sem->Type()) : nullptr;
-}
-
-std::string Resolver::TypeNameOf(const sem::Type* ty) {
-  return RawTypeNameOf(ty->UnwrapRef());
-}
-
-std::string Resolver::RawTypeNameOf(const sem::Type* ty) {
-  return ty->FriendlyName(builder_->Symbols());
-}
-
-sem::Type* Resolver::TypeOf(const ast::LiteralExpression* lit) {
-  return Switch(
-      lit,
-      [&](const ast::SintLiteralExpression*) {
-        return builder_->create<sem::I32>();
-      },
-      [&](const ast::UintLiteralExpression*) {
-        return builder_->create<sem::U32>();
-      },
-      [&](const ast::FloatLiteralExpression*) {
-        return builder_->create<sem::F32>();
-      },
-      [&](const ast::BoolLiteralExpression*) {
-        return builder_->create<sem::Bool>();
-      },
-      [&](Default) {
-        TINT_UNREACHABLE(Resolver, diagnostics_)
-            << "Unhandled literal type: " << lit->TypeInfo().name;
-        return nullptr;
-      });
-}
-
 sem::Array* Resolver::Array(const ast::Array* arr) {
   auto source = arr->source;
 
@@ -2184,7 +2157,7 @@
   }
 
   if (!IsPlain(elem_type)) {  // Check must come before GetDefaultAlignAndSize()
-    AddError(TypeNameOf(elem_type) +
+    AddError(sem_.TypeNameOf(elem_type) +
                  " cannot be used as an element type of an array",
              source);
     return nullptr;
@@ -2364,7 +2337,7 @@
 
     // Validate member type
     if (!IsPlain(type)) {
-      AddError(TypeNameOf(type) +
+      AddError(sem_.TypeNameOf(type) +
                    " cannot be used as the type of a structure member",
                member->source);
       return nullptr;
@@ -2477,7 +2450,10 @@
     }
   }
 
-  if (!ValidateStructure(out)) {
+  auto stage = current_function_
+                   ? current_function_->Declaration()->PipelineStage()
+                   : ast::PipelineStage::kNone;
+  if (!ValidateStructure(out, stage)) {
     return nullptr;
   }
 
@@ -2501,7 +2477,9 @@
 
     // Validate after processing the return value expression so that its type
     // is available for validation.
-    return ValidateReturn(stmt);
+    auto* ret_type = stmt->value ? sem_.TypeOf(stmt->value)->UnwrapRef()
+                                 : builder_->create<sem::Void>();
+    return ValidateReturn(stmt, current_function_->ReturnType(), ret_type);
   });
 }
 
@@ -2589,7 +2567,7 @@
       behaviors.Add(lhs->Behaviors());
     }
 
-    return ValidateAssignment(stmt, TypeOf(stmt->rhs));
+    return ValidateAssignment(stmt, sem_.TypeOf(stmt->rhs));
   });
 }
 
@@ -2637,8 +2615,8 @@
     auto* ty = BinaryOpType(lhs_ty, rhs_ty, stmt->op);
     if (!ty) {
       AddError("compound assignment operand types are invalid: " +
-                   TypeNameOf(lhs_ty) + " " + FriendlyName(stmt->op) + " " +
-                   TypeNameOf(rhs_ty),
+                   sem_.TypeNameOf(lhs_ty) + " " + FriendlyName(stmt->op) +
+                   " " + sem_.TypeNameOf(rhs_ty),
                stmt->source);
       return false;
     }
@@ -2717,7 +2695,8 @@
     for (auto* member : str->Members()) {
       if (!ApplyStorageClassUsageToType(sc, member->Type(), usage)) {
         std::stringstream err;
-        err << "while analysing structure member " << TypeNameOf(str) << "."
+        err << "while analysing structure member " << sem_.TypeNameOf(str)
+            << "."
             << builder_->Symbols().NameFor(member->Declaration()->symbol);
         AddNote(err.str(), member->Declaration()->source);
         return false;
@@ -2741,8 +2720,9 @@
 
   if (ast::IsHostShareable(sc) && !IsHostShareable(ty)) {
     std::stringstream err;
-    err << "Type '" << TypeNameOf(ty) << "' cannot be used in storage class '"
-        << sc << "' as it is non-host-shareable";
+    err << "Type '" << sem_.TypeNameOf(ty)
+        << "' cannot be used in storage class '" << sc
+        << "' as it is non-host-shareable";
     AddError(err.str(), usage);
     return false;
   }
@@ -2774,12 +2754,6 @@
   return sem;
 }
 
-std::string Resolver::VectorPretty(uint32_t size,
-                                   const sem::Type* element_type) {
-  sem::Vector vec_type(element_type, size);
-  return vec_type.FriendlyName(builder_->Symbols());
-}
-
 bool Resolver::Mark(const ast::Node* node) {
   if (node == nullptr) {
     TINT_ICE(Resolver, diagnostics_) << "Resolver::Mark() called with nullptr";
diff --git a/src/tint/resolver/resolver.h b/src/tint/resolver/resolver.h
index ff879b0..95dd5ff 100644
--- a/src/tint/resolver/resolver.h
+++ b/src/tint/resolver/resolver.h
@@ -26,6 +26,7 @@
 #include "src/tint/builtin_table.h"
 #include "src/tint/program_builder.h"
 #include "src/tint/resolver/dependency_graph.h"
+#include "src/tint/resolver/sem_helper.h"
 #include "src/tint/scope_stack.h"
 #include "src/tint/sem/binding_point.h"
 #include "src/tint/sem/block_statement.h"
@@ -106,8 +107,9 @@
   /// Describes the context in which a variable is declared
   enum class VariableKind { kParameter, kLocal, kGlobal };
 
-  std::set<std::pair<const sem::Type*, ast::StorageClass>>
-      valid_type_storage_layouts_;
+  using ValidTypeStorageLayouts =
+      std::set<std::pair<const sem::Type*, ast::StorageClass>>;
+  ValidTypeStorageLayouts valid_type_storage_layouts_;
 
   /// Structure holding semantic information about a block (i.e. scope), such as
   /// parent block and variables declared in the block.
@@ -153,8 +155,6 @@
   /// @returns true on success, false on error
   bool ResolveInternal();
 
-  bool ValidatePipelineStages();
-
   /// Creates the nodes and adds them to the sem::Info mappings of the
   /// ProgramBuilder.
   void CreateSemanticNodes() const;
@@ -239,72 +239,84 @@
 
   // AST and Type validation methods
   // Each return true on success, false on failure.
-  bool ValidateAlias(const ast::Alias*);
-  bool ValidateArray(const sem::Array* arr, const Source& source);
+  bool ValidatePipelineStages() const;
+  bool ValidateAlias(const ast::Alias*) const;
+  bool ValidateArray(const sem::Array* arr, const Source& source) const;
   bool ValidateArrayStrideAttribute(const ast::StrideAttribute* attr,
                                     uint32_t el_size,
                                     uint32_t el_align,
-                                    const Source& source);
-  bool ValidateAtomic(const ast::Atomic* a, const sem::Atomic* s);
-  bool ValidateAtomicVariable(const sem::Variable* var);
-  bool ValidateAssignment(const ast::Statement* a, const sem::Type* rhs_ty);
-  bool ValidateBitcast(const ast::BitcastExpression* cast, const sem::Type* to);
-  bool ValidateBreakStatement(const sem::Statement* stmt);
+                                    const Source& source) const;
+  bool ValidateAtomic(const ast::Atomic* a, const sem::Atomic* s) const;
+  bool ValidateAtomicVariable(const sem::Variable* var) const;
+  bool ValidateAssignment(const ast::Statement* a,
+                          const sem::Type* rhs_ty) const;
+  bool ValidateBitcast(const ast::BitcastExpression* cast,
+                       const sem::Type* to) const;
+  bool ValidateBreakStatement(const sem::Statement* stmt) const;
   bool ValidateBuiltinAttribute(const ast::BuiltinAttribute* attr,
                                 const sem::Type* storage_type,
-                                const bool is_input);
-  bool ValidateContinueStatement(const sem::Statement* stmt);
-  bool ValidateDiscardStatement(const sem::Statement* stmt);
-  bool ValidateElseStatement(const sem::ElseStatement* stmt);
-  bool ValidateEntryPoint(const sem::Function* func);
-  bool ValidateForLoopStatement(const sem::ForLoopStatement* stmt);
-  bool ValidateFallthroughStatement(const sem::Statement* stmt);
-  bool ValidateFunction(const sem::Function* func);
-  bool ValidateFunctionCall(const sem::Call* call);
-  bool ValidateGlobalVariable(const sem::Variable* var);
-  bool ValidateIfStatement(const sem::IfStatement* stmt);
+                                ast::PipelineStage stage,
+                                const bool is_input) const;
+  bool ValidateContinueStatement(const sem::Statement* stmt) const;
+  bool ValidateDiscardStatement(const sem::Statement* stmt) const;
+  bool ValidateElseStatement(const sem::ElseStatement* stmt) const;
+  bool ValidateEntryPoint(const sem::Function* func,
+                          ast::PipelineStage stage) const;
+  bool ValidateForLoopStatement(const sem::ForLoopStatement* stmt) const;
+  bool ValidateFallthroughStatement(const sem::Statement* stmt) const;
+  bool ValidateFunction(const sem::Function* func,
+                        ast::PipelineStage stage) const;
+  bool ValidateFunctionCall(const sem::Call* call) const;
+  bool ValidateGlobalVariable(const sem::Variable* var) const;
+  bool ValidateIfStatement(const sem::IfStatement* stmt) const;
   bool ValidateIncrementDecrementStatement(
-      const ast::IncrementDecrementStatement* stmt);
+      const ast::IncrementDecrementStatement* stmt) const;
   bool ValidateInterpolateAttribute(const ast::InterpolateAttribute* attr,
-                                    const sem::Type* storage_type);
-  bool ValidateBuiltinCall(const sem::Call* call);
+                                    const sem::Type* storage_type) const;
+  bool ValidateBuiltinCall(const sem::Call* call) const;
   bool ValidateLocationAttribute(const ast::LocationAttribute* location,
                                  const sem::Type* type,
                                  std::unordered_set<uint32_t>& locations,
+                                 ast::PipelineStage stage,
                                  const Source& source,
-                                 const bool is_input = false);
-  bool ValidateLoopStatement(const sem::LoopStatement* stmt);
-  bool ValidateMatrix(const sem::Matrix* ty, const Source& source);
+                                 const bool is_input = false) const;
+  bool ValidateLoopStatement(const sem::LoopStatement* stmt) const;
+  bool ValidateMatrix(const sem::Matrix* ty, const Source& source) const;
   bool ValidateFunctionParameter(const ast::Function* func,
-                                 const sem::Variable* var);
-  bool ValidateParameter(const ast::Function* func, const sem::Variable* var);
-  bool ValidateReturn(const ast::ReturnStatement* ret);
-  bool ValidateStatements(const ast::StatementList& stmts);
-  bool ValidateStorageTexture(const ast::StorageTexture* t);
-  bool ValidateStructure(const sem::Struct* str);
+                                 const sem::Variable* var) const;
+  bool ValidateReturn(const ast::ReturnStatement* ret,
+                      const sem::Type* func_type,
+                      const sem::Type* ret_type) const;
+  bool ValidateStatements(const ast::StatementList& stmts) const;
+  bool ValidateStorageTexture(const ast::StorageTexture* t) const;
+  bool ValidateStructure(const sem::Struct* str,
+                         ast::PipelineStage stage) const;
   bool ValidateStructureConstructorOrCast(const ast::CallExpression* ctor,
-                                          const sem::Struct* struct_type);
+                                          const sem::Struct* struct_type) const;
   bool ValidateSwitch(const ast::SwitchStatement* s);
-  bool ValidateVariable(const sem::Variable* var);
+  bool ValidateVariable(const sem::Variable* var) const;
   bool ValidateVariableConstructorOrCast(const ast::Variable* var,
                                          ast::StorageClass storage_class,
                                          const sem::Type* storage_type,
-                                         const sem::Type* rhs_type);
-  bool ValidateVector(const sem::Vector* ty, const Source& source);
+                                         const sem::Type* rhs_type) const;
+  bool ValidateVector(const sem::Vector* ty, const Source& source) const;
   bool ValidateVectorConstructorOrCast(const ast::CallExpression* ctor,
-                                       const sem::Vector* vec_type);
+                                       const sem::Vector* vec_type) const;
   bool ValidateMatrixConstructorOrCast(const ast::CallExpression* ctor,
-                                       const sem::Matrix* matrix_type);
+                                       const sem::Matrix* matrix_type) const;
   bool ValidateScalarConstructorOrCast(const ast::CallExpression* ctor,
-                                       const sem::Type* type);
+                                       const sem::Type* type) const;
   bool ValidateArrayConstructorOrCast(const ast::CallExpression* ctor,
-                                      const sem::Array* arr_type);
-  bool ValidateTextureBuiltinFunction(const sem::Call* call);
-  bool ValidateNoDuplicateAttributes(const ast::AttributeList& attributes);
+                                      const sem::Array* arr_type) const;
+  bool ValidateTextureBuiltinFunction(const sem::Call* call) const;
+  bool ValidateNoDuplicateAttributes(
+      const ast::AttributeList& attributes) const;
   bool ValidateStorageClassLayout(const sem::Type* type,
                                   ast::StorageClass sc,
-                                  Source source);
-  bool ValidateStorageClassLayout(const sem::Variable* var);
+                                  Source source,
+                                  ValidTypeStorageLayouts& layouts) const;
+  bool ValidateStorageClassLayout(const sem::Variable* var,
+                                  ValidTypeStorageLayouts& layouts) const;
 
   /// @returns true if the attribute list contains a
   /// ast::DisableValidationAttribute with the validation mode equal to
@@ -318,6 +330,13 @@
   bool IsValidationEnabled(const ast::AttributeList& attributes,
                            ast::DisabledValidation validation) const;
 
+  /// Returns a human-readable string representation of the vector type name
+  /// with the given parameters.
+  /// @param size the vector dimension
+  /// @param element_type scalar vector sub-element type
+  /// @return pretty string representation
+  std::string VectorPretty(uint32_t size, const sem::Type* element_type) const;
+
   /// Resolves the WorkgroupSize for the given function, assigning it to
   /// current_function_
   bool WorkgroupSize(const ast::Function*);
@@ -387,23 +406,6 @@
   /// Set the shadowing information on variable declarations.
   /// @note this method must only be called after all semantic nodes are built.
   void SetShadows();
-
-  /// @returns the resolved type of the ast::Expression `expr`
-  /// @param expr the expression
-  sem::Type* TypeOf(const ast::Expression* expr);
-
-  /// @returns the type name of the given semantic type, unwrapping
-  /// references.
-  std::string TypeNameOf(const sem::Type* ty);
-
-  /// @returns the type name of the given semantic type, without unwrapping
-  /// references.
-  std::string RawTypeNameOf(const sem::Type* ty);
-
-  /// @returns the semantic type of the AST literal `lit`
-  /// @param lit the literal
-  sem::Type* TypeOf(const ast::LiteralExpression* lit);
-
   /// StatementScope() does the following:
   /// * Creates the AST -> SEM mapping.
   /// * Assigns `sem` to #current_statement_
@@ -418,13 +420,6 @@
   template <typename SEM, typename F>
   SEM* StatementScope(const ast::Statement* ast, SEM* sem, F&& callback);
 
-  /// Returns a human-readable string representation of the vector type name
-  /// with the given parameters.
-  /// @param size the vector dimension
-  /// @param element_type scalar vector sub-element type
-  /// @return pretty string representation
-  std::string VectorPretty(uint32_t size, const sem::Type* element_type);
-
   /// Mark records that the given AST node has been visited, and asserts that
   /// the given node has not already been seen. Diamonds in the AST are
   /// illegal.
@@ -456,21 +451,6 @@
   sem::Constant EvaluateConstantValue(const ast::CallExpression* call,
                                       const sem::Type* type);
 
-  /// Sem is a helper for obtaining the semantic node for the given AST node.
-  template <typename SEM = sem::Info::InferFromAST,
-            typename AST_OR_TYPE = CastableBase>
-  auto* Sem(const AST_OR_TYPE* ast) {
-    using T = sem::Info::GetResultType<SEM, AST_OR_TYPE>;
-    auto* sem = builder_->Sem().Get(ast);
-    if (!sem) {
-      TINT_ICE(Resolver, diagnostics_)
-          << "AST node '" << ast->TypeInfo().name << "' had no semantic info\n"
-          << "At: " << ast->source << "\n"
-          << "Pointer: " << ast;
-    }
-    return const_cast<T*>(As<T>(sem));
-  }
-
   /// @returns true if the symbol is the name of a builtin function.
   bool IsBuiltin(Symbol) const;
 
@@ -488,7 +468,7 @@
   /// @returns the resolved symbol (function, type or variable) for the given
   /// ast::Identifier or ast::TypeName cast to the given semantic type.
   template <typename SEM = sem::Node>
-  SEM* ResolvedSymbol(const ast::Node* node) {
+  SEM* ResolvedSymbol(const ast::Node* node) const {
     auto* resolved = utils::Lookup(dependencies_.resolved_symbols, node);
     return resolved ? const_cast<SEM*>(builder_->Sem().Get<SEM>(resolved))
                     : nullptr;
@@ -530,6 +510,7 @@
   diag::List& diagnostics_;
   std::unique_ptr<BuiltinTable> const builtin_table_;
   DependencyGraph dependencies_;
+  SemHelper sem_;
   std::vector<sem::Function*> entry_points_;
   std::unordered_map<const sem::Type*, const Source&> atomic_composite_info_;
   std::unordered_set<const ast::Node*> marked_;
diff --git a/src/tint/resolver/resolver_validation.cc b/src/tint/resolver/resolver_validation.cc
index 708a4fa..44e0a2f 100644
--- a/src/tint/resolver/resolver_validation.cc
+++ b/src/tint/resolver/resolver_validation.cc
@@ -149,7 +149,8 @@
 
 }  // namespace
 
-bool Resolver::ValidateAtomic(const ast::Atomic* a, const sem::Atomic* s) {
+bool Resolver::ValidateAtomic(const ast::Atomic* a,
+                              const sem::Atomic* s) const {
   // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types
   // T must be either u32 or i32.
   if (!s->Type()->IsAnyOf<sem::U32, sem::I32>()) {
@@ -160,7 +161,7 @@
   return true;
 }
 
-bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) {
+bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) const {
   switch (t->access) {
     case ast::Access::kWrite:
       break;
@@ -193,15 +194,15 @@
     const ast::Variable* var,
     ast::StorageClass storage_class,
     const sem::Type* storage_ty,
-    const sem::Type* rhs_ty) {
+    const sem::Type* rhs_ty) const {
   auto* value_type = rhs_ty->UnwrapRef();  // Implicit load of RHS
 
   // Value type has to match storage type
   if (storage_ty != value_type) {
     std::string decl = var->is_const ? "let" : "var";
     AddError("cannot initialize " + decl + " of type '" +
-                 TypeNameOf(storage_ty) + "' with value of type '" +
-                 TypeNameOf(rhs_ty) + "'",
+                 sem_.TypeNameOf(storage_ty) + "' with value of type '" +
+                 sem_.TypeNameOf(rhs_ty) + "'",
              var->source);
     return false;
   }
@@ -228,9 +229,11 @@
   return true;
 }
 
-bool Resolver::ValidateStorageClassLayout(const sem::Type* store_ty,
-                                          ast::StorageClass sc,
-                                          Source source) {
+bool Resolver::ValidateStorageClassLayout(
+    const sem::Type* store_ty,
+    ast::StorageClass sc,
+    Source source,
+    ValidTypeStorageLayouts& layouts) const {
   // https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints
 
   auto is_uniform_struct_or_array = [sc](const sem::Type* ty) {
@@ -256,7 +259,7 @@
   };
 
   // Cache result of type + storage class pair.
-  if (!valid_type_storage_layouts_.emplace(store_ty, sc).second) {
+  if (!layouts.emplace(store_ty, sc).second) {
     return true;
   }
 
@@ -270,8 +273,8 @@
       uint32_t required_align = required_alignment_of(m->Type());
 
       // Recurse into the member type.
-      if (!ValidateStorageClassLayout(m->Type(), sc,
-                                      m->Declaration()->type->source)) {
+      if (!ValidateStorageClassLayout(
+              m->Type(), sc, m->Declaration()->type->source, layouts)) {
         AddNote("see layout of struct:\n" + str->Layout(builder_->Symbols()),
                 str->Declaration()->source);
         return false;
@@ -339,7 +342,7 @@
     // TODO(crbug.com/tint/1388): Ideally we'd pass the source for nested
     // element type here, but we can't easily get that from the semantic node.
     // We should consider recursing through the AST type nodes instead.
-    if (!ValidateStorageClassLayout(arr->ElemType(), sc, source)) {
+    if (!ValidateStorageClassLayout(arr->ElemType(), sc, source, layouts)) {
       return false;
     }
 
@@ -381,10 +384,12 @@
   return true;
 }
 
-bool Resolver::ValidateStorageClassLayout(const sem::Variable* var) {
+bool Resolver::ValidateStorageClassLayout(
+    const sem::Variable* var,
+    ValidTypeStorageLayouts& layouts) const {
   if (auto* str = var->Type()->UnwrapRef()->As<sem::Struct>()) {
     if (!ValidateStorageClassLayout(str, var->StorageClass(),
-                                    str->Declaration()->source)) {
+                                    str->Declaration()->source, layouts)) {
       AddNote("see declaration of variable", var->Declaration()->source);
       return false;
     }
@@ -394,7 +399,7 @@
       source = var->Declaration()->type->source;
     }
     if (!ValidateStorageClassLayout(var->Type()->UnwrapRef(),
-                                    var->StorageClass(), source)) {
+                                    var->StorageClass(), source, layouts)) {
       return false;
     }
   }
@@ -402,7 +407,7 @@
   return true;
 }
 
-bool Resolver::ValidateGlobalVariable(const sem::Variable* var) {
+bool Resolver::ValidateGlobalVariable(const sem::Variable* var) const {
   auto* decl = var->Declaration();
   if (!ValidateNoDuplicateAttributes(decl->attributes)) {
     return false;
@@ -508,7 +513,7 @@
 // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types
 // Atomic types may only be instantiated by variables in the workgroup storage
 // class or by storage buffer variables with a read_write access mode.
-bool Resolver::ValidateAtomicVariable(const sem::Variable* var) {
+bool Resolver::ValidateAtomicVariable(const sem::Variable* var) const {
   auto sc = var->StorageClass();
   auto* decl = var->Declaration();
   auto access = var->Access();
@@ -531,9 +536,9 @@
         AddError(
             "atomic variables must have <storage> or <workgroup> storage class",
             source);
-        AddNote(
-            "atomic sub-type of '" + TypeNameOf(type) + "' is declared here",
-            found->second);
+        AddNote("atomic sub-type of '" + sem_.TypeNameOf(type) +
+                    "' is declared here",
+                found->second);
         return false;
       } else if (sc == ast::StorageClass::kStorage &&
                  access != ast::Access::kReadWrite) {
@@ -541,9 +546,9 @@
             "atomic variables in <storage> storage class must have read_write "
             "access mode",
             source);
-        AddNote(
-            "atomic sub-type of '" + TypeNameOf(type) + "' is declared here",
-            found->second);
+        AddNote("atomic sub-type of '" + sem_.TypeNameOf(type) +
+                    "' is declared here",
+                found->second);
         return false;
       }
     }
@@ -552,7 +557,7 @@
   return true;
 }
 
-bool Resolver::ValidateVariable(const sem::Variable* var) {
+bool Resolver::ValidateVariable(const sem::Variable* var) const {
   auto* decl = var->Declaration();
   auto* storage_ty = var->Type()->UnwrapRef();
 
@@ -570,15 +575,17 @@
   }
 
   if (!decl->is_const && !IsStorable(storage_ty)) {
-    AddError(TypeNameOf(storage_ty) + " cannot be used as the type of a var",
-             decl->source);
+    AddError(
+        sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a var",
+        decl->source);
     return false;
   }
 
   if (decl->is_const && !var->Is<sem::Parameter>() &&
       !(storage_ty->IsConstructible() || storage_ty->Is<sem::Pointer>())) {
-    AddError(TypeNameOf(storage_ty) + " cannot be used as the type of a let",
-             decl->source);
+    AddError(
+        sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a let",
+        decl->source);
     return false;
   }
 
@@ -611,7 +618,7 @@
     // If the store type is a texture type or a sampler type, then the
     // variable declaration must not have a storage class attribute. The
     // storage class will always be handle.
-    AddError("variables of type '" + TypeNameOf(storage_ty) +
+    AddError("variables of type '" + sem_.TypeNameOf(storage_ty) +
                  "' must not have a storage class",
              decl->source);
     return false;
@@ -628,7 +635,7 @@
 }
 
 bool Resolver::ValidateFunctionParameter(const ast::Function* func,
-                                         const sem::Variable* var) {
+                                         const sem::Variable* var) const {
   if (!ValidateVariable(var)) {
     return false;
   }
@@ -681,9 +688,9 @@
     }
   } else if (!var->Type()
                   ->IsAnyOf<sem::Texture, sem::Sampler, sem::Pointer>()) {
-    AddError(
-        "store type of function parameter cannot be " + TypeNameOf(var->Type()),
-        decl->source);
+    AddError("store type of function parameter cannot be " +
+                 sem_.TypeNameOf(var->Type()),
+             decl->source);
     return false;
   }
 
@@ -692,11 +699,9 @@
 
 bool Resolver::ValidateBuiltinAttribute(const ast::BuiltinAttribute* attr,
                                         const sem::Type* storage_ty,
-                                        const bool is_input) {
+                                        ast::PipelineStage stage,
+                                        const bool is_input) const {
   auto* type = storage_ty->UnwrapRef();
-  const auto stage = current_function_
-                         ? current_function_->Declaration()->PipelineStage()
-                         : ast::PipelineStage::kNone;
   std::stringstream stage_name;
   stage_name << stage;
   bool is_stage_mismatch = false;
@@ -813,7 +818,7 @@
 
 bool Resolver::ValidateInterpolateAttribute(
     const ast::InterpolateAttribute* attr,
-    const sem::Type* storage_ty) {
+    const sem::Type* storage_ty) const {
   auto* type = storage_ty->UnwrapRef();
 
   if (type->is_integer_scalar_or_vector() &&
@@ -834,7 +839,8 @@
   return true;
 }
 
-bool Resolver::ValidateFunction(const sem::Function* func) {
+bool Resolver::ValidateFunction(const sem::Function* func,
+                                ast::PipelineStage stage) const {
   auto* decl = func->Declaration();
 
   auto name = builder_->Symbols().NameFor(decl->symbol);
@@ -882,7 +888,7 @@
     if (decl->body) {
       sem::Behaviors behaviors{sem::Behavior::kNext};
       if (auto* last = decl->body->Last()) {
-        behaviors = Sem(last)->Behaviors();
+        behaviors = sem_.Get(last)->Behaviors();
       }
       if (behaviors.Contains(sem::Behavior::kNext)) {
         AddError("missing return at end of function", decl->source);
@@ -919,7 +925,7 @@
   }
 
   if (decl->IsEntryPoint()) {
-    if (!ValidateEntryPoint(func)) {
+    if (!ValidateEntryPoint(func, stage)) {
       return false;
     }
   }
@@ -939,7 +945,8 @@
   return true;
 }
 
-bool Resolver::ValidateEntryPoint(const sem::Function* func) {
+bool Resolver::ValidateEntryPoint(const sem::Function* func,
+                                  ast::PipelineStage stage) const {
   auto* decl = func->Declaration();
 
   // Use a lambda to validate the entry point attributes for a type.
@@ -968,6 +975,7 @@
     const ast::InvariantAttribute* invariant_attribute = nullptr;
     for (auto* attr : attrs) {
       auto is_invalid_compute_shader_attribute = false;
+
       if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
         if (pipeline_io_attribute) {
           AddError("multiple entry point IO attributes", attr->source);
@@ -987,7 +995,7 @@
         }
 
         if (!ValidateBuiltinAttribute(
-                builtin, ty,
+                builtin, ty, stage,
                 /* is_input */ param_or_ret == ParamOrRetType::kParameter)) {
           return false;
         }
@@ -1002,7 +1010,8 @@
         pipeline_io_attribute = attr;
 
         bool is_input = param_or_ret == ParamOrRetType::kParameter;
-        if (!ValidateLocationAttribute(location, ty, locations, source,
+
+        if (!ValidateLocationAttribute(location, ty, locations, stage, source,
                                        is_input)) {
           return false;
         }
@@ -1213,9 +1222,9 @@
   return true;
 }
 
-bool Resolver::ValidateStatements(const ast::StatementList& stmts) {
+bool Resolver::ValidateStatements(const ast::StatementList& stmts) const {
   for (auto* stmt : stmts) {
-    if (!Sem(stmt)->IsReachable()) {
+    if (!sem_.Get(stmt)->IsReachable()) {
       /// TODO(https://github.com/gpuweb/gpuweb/issues/2378): This may need to
       /// become an error.
       AddWarning("code is unreachable", stmt->source);
@@ -1226,15 +1235,16 @@
 }
 
 bool Resolver::ValidateBitcast(const ast::BitcastExpression* cast,
-                               const sem::Type* to) {
-  auto* from = TypeOf(cast->expr)->UnwrapRef();
+                               const sem::Type* to) const {
+  auto* from = sem_.TypeOf(cast->expr)->UnwrapRef();
   if (!from->is_numeric_scalar_or_vector()) {
-    AddError("'" + TypeNameOf(from) + "' cannot be bitcast",
+    AddError("'" + sem_.TypeNameOf(from) + "' cannot be bitcast",
              cast->expr->source);
     return false;
   }
   if (!to->is_numeric_scalar_or_vector()) {
-    AddError("cannot bitcast to '" + TypeNameOf(to) + "'", cast->type->source);
+    AddError("cannot bitcast to '" + sem_.TypeNameOf(to) + "'",
+             cast->type->source);
     return false;
   }
 
@@ -1246,8 +1256,8 @@
   };
 
   if (width(from) != width(to)) {
-    AddError("cannot bitcast from '" + TypeNameOf(from) + "' to '" +
-                 TypeNameOf(to) + "'",
+    AddError("cannot bitcast from '" + sem_.TypeNameOf(from) + "' to '" +
+                 sem_.TypeNameOf(to) + "'",
              cast->source);
     return false;
   }
@@ -1255,7 +1265,7 @@
   return true;
 }
 
-bool Resolver::ValidateBreakStatement(const sem::Statement* stmt) {
+bool Resolver::ValidateBreakStatement(const sem::Statement* stmt) const {
   if (!stmt->FindFirstParent<sem::LoopBlockStatement, sem::CaseStatement>()) {
     AddError("break statement must be in a loop or switch case",
              stmt->Declaration()->source);
@@ -1322,7 +1332,7 @@
   return true;
 }
 
-bool Resolver::ValidateContinueStatement(const sem::Statement* stmt) {
+bool Resolver::ValidateContinueStatement(const sem::Statement* stmt) const {
   if (auto* continuing = ClosestContinuing(/*stop_at_loop*/ true)) {
     AddError("continuing blocks must not contain a continue statement",
              stmt->Declaration()->source);
@@ -1342,7 +1352,7 @@
   return true;
 }
 
-bool Resolver::ValidateDiscardStatement(const sem::Statement* stmt) {
+bool Resolver::ValidateDiscardStatement(const sem::Statement* stmt) const {
   if (auto* continuing = ClosestContinuing(/*stop_at_loop*/ false)) {
     AddError("continuing blocks must not contain a discard statement",
              stmt->Declaration()->source);
@@ -1355,7 +1365,7 @@
   return true;
 }
 
-bool Resolver::ValidateFallthroughStatement(const sem::Statement* stmt) {
+bool Resolver::ValidateFallthroughStatement(const sem::Statement* stmt) const {
   if (auto* block = As<sem::BlockStatement>(stmt->Parent())) {
     if (auto* c = As<sem::CaseStatement>(block->Parent())) {
       if (block->Declaration()->Last() == stmt->Declaration()) {
@@ -1378,36 +1388,12 @@
   return false;
 }
 
-bool Resolver::ValidateElseStatement(const sem::ElseStatement* stmt) {
+bool Resolver::ValidateElseStatement(const sem::ElseStatement* stmt) const {
   if (auto* cond = stmt->Condition()) {
     auto* cond_ty = cond->Type()->UnwrapRef();
     if (!cond_ty->Is<sem::Bool>()) {
-      AddError(
-          "else statement condition must be bool, got " + TypeNameOf(cond_ty),
-          stmt->Condition()->Declaration()->source);
-      return false;
-    }
-  }
-  return true;
-}
-
-bool Resolver::ValidateLoopStatement(const sem::LoopStatement* stmt) {
-  if (stmt->Behaviors().Empty()) {
-    AddError("loop does not exit", stmt->Declaration()->source.Begin());
-    return false;
-  }
-  return true;
-}
-
-bool Resolver::ValidateForLoopStatement(const sem::ForLoopStatement* stmt) {
-  if (stmt->Behaviors().Empty()) {
-    AddError("for-loop does not exit", stmt->Declaration()->source.Begin());
-    return false;
-  }
-  if (auto* cond = stmt->Condition()) {
-    auto* cond_ty = cond->Type()->UnwrapRef();
-    if (!cond_ty->Is<sem::Bool>()) {
-      AddError("for-loop condition must be bool, got " + TypeNameOf(cond_ty),
+      AddError("else statement condition must be bool, got " +
+                   sem_.TypeNameOf(cond_ty),
                stmt->Condition()->Declaration()->source);
       return false;
     }
@@ -1415,17 +1401,44 @@
   return true;
 }
 
-bool Resolver::ValidateIfStatement(const sem::IfStatement* stmt) {
-  auto* cond_ty = stmt->Condition()->Type()->UnwrapRef();
-  if (!cond_ty->Is<sem::Bool>()) {
-    AddError("if statement condition must be bool, got " + TypeNameOf(cond_ty),
-             stmt->Condition()->Declaration()->source);
+bool Resolver::ValidateLoopStatement(const sem::LoopStatement* stmt) const {
+  if (stmt->Behaviors().Empty()) {
+    AddError("loop does not exit", stmt->Declaration()->source.Begin());
     return false;
   }
   return true;
 }
 
-bool Resolver::ValidateBuiltinCall(const sem::Call* call) {
+bool Resolver::ValidateForLoopStatement(
+    const sem::ForLoopStatement* stmt) const {
+  if (stmt->Behaviors().Empty()) {
+    AddError("for-loop does not exit", stmt->Declaration()->source.Begin());
+    return false;
+  }
+  if (auto* cond = stmt->Condition()) {
+    auto* cond_ty = cond->Type()->UnwrapRef();
+    if (!cond_ty->Is<sem::Bool>()) {
+      AddError(
+          "for-loop condition must be bool, got " + sem_.TypeNameOf(cond_ty),
+          stmt->Condition()->Declaration()->source);
+      return false;
+    }
+  }
+  return true;
+}
+
+bool Resolver::ValidateIfStatement(const sem::IfStatement* stmt) const {
+  auto* cond_ty = stmt->Condition()->Type()->UnwrapRef();
+  if (!cond_ty->Is<sem::Bool>()) {
+    AddError(
+        "if statement condition must be bool, got " + sem_.TypeNameOf(cond_ty),
+        stmt->Condition()->Declaration()->source);
+    return false;
+  }
+  return true;
+}
+
+bool Resolver::ValidateBuiltinCall(const sem::Call* call) const {
   if (call->Type()->Is<sem::Void>()) {
     bool is_call_statement = false;
     if (auto* call_stmt = As<ast::CallStatement>(call->Stmt()->Declaration())) {
@@ -1448,7 +1461,7 @@
   return true;
 }
 
-bool Resolver::ValidateTextureBuiltinFunction(const sem::Call* call) {
+bool Resolver::ValidateTextureBuiltinFunction(const sem::Call* call) const {
   auto* builtin = call->Target()->As<sem::Builtin>();
   if (!builtin) {
     return false;
@@ -1520,7 +1533,7 @@
          check_arg_is_constexpr(sem::ParameterUsage::kComponent, 0, 3);
 }
 
-bool Resolver::ValidateFunctionCall(const sem::Call* call) {
+bool Resolver::ValidateFunctionCall(const sem::Call* call) const {
   auto* decl = call->Declaration();
   auto* target = call->Target()->As<sem::Function>();
   auto sym = decl->target.name->symbol;
@@ -1548,13 +1561,13 @@
     const sem::Variable* param = target->Parameters()[i];
     const ast::Expression* arg_expr = decl->args[i];
     auto* param_type = param->Type();
-    auto* arg_type = TypeOf(arg_expr)->UnwrapRef();
+    auto* arg_type = sem_.TypeOf(arg_expr)->UnwrapRef();
 
     if (param_type != arg_type) {
       AddError("type mismatch for argument " + std::to_string(i + 1) +
                    " in call to '" + name + "', expected '" +
-                   TypeNameOf(param_type) + "', got '" + TypeNameOf(arg_type) +
-                   "'",
+                   sem_.TypeNameOf(param_type) + "', got '" +
+                   sem_.TypeNameOf(arg_type) + "'",
                arg_expr->source);
       return false;
     }
@@ -1638,7 +1651,7 @@
 
 bool Resolver::ValidateStructureConstructorOrCast(
     const ast::CallExpression* ctor,
-    const sem::Struct* struct_type) {
+    const sem::Struct* struct_type) const {
   if (!struct_type->IsConstructible()) {
     AddError("struct constructor has non-constructible type", ctor->source);
     return false;
@@ -1656,13 +1669,13 @@
     }
     for (auto* member : struct_type->Members()) {
       auto* value = ctor->args[member->Index()];
-      auto* value_ty = TypeOf(value);
+      auto* value_ty = sem_.TypeOf(value);
       if (member->Type() != value_ty->UnwrapRef()) {
         AddError(
             "type in struct constructor does not match struct member type: "
             "expected '" +
-                TypeNameOf(member->Type()) + "', found '" +
-                TypeNameOf(value_ty) + "'",
+                sem_.TypeNameOf(member->Type()) + "', found '" +
+                sem_.TypeNameOf(value_ty) + "'",
             value->source);
         return false;
       }
@@ -1671,17 +1684,19 @@
   return true;
 }
 
-bool Resolver::ValidateArrayConstructorOrCast(const ast::CallExpression* ctor,
-                                              const sem::Array* array_type) {
+bool Resolver::ValidateArrayConstructorOrCast(
+    const ast::CallExpression* ctor,
+    const sem::Array* array_type) const {
   auto& values = ctor->args;
   auto* elem_ty = array_type->ElemType();
   for (auto* value : values) {
-    auto* value_ty = TypeOf(value)->UnwrapRef();
+    auto* value_ty = sem_.TypeOf(value)->UnwrapRef();
     if (value_ty != elem_ty) {
       AddError(
           "type in array constructor does not match array type: "
           "expected '" +
-              TypeNameOf(elem_ty) + "', found '" + TypeNameOf(value_ty) + "'",
+              sem_.TypeNameOf(elem_ty) + "', found '" +
+              sem_.TypeNameOf(value_ty) + "'",
           value->source);
       return false;
     }
@@ -1711,19 +1726,21 @@
   return true;
 }
 
-bool Resolver::ValidateVectorConstructorOrCast(const ast::CallExpression* ctor,
-                                               const sem::Vector* vec_type) {
+bool Resolver::ValidateVectorConstructorOrCast(
+    const ast::CallExpression* ctor,
+    const sem::Vector* vec_type) const {
   auto& values = ctor->args;
   auto* elem_ty = vec_type->type();
   size_t value_cardinality_sum = 0;
   for (auto* value : values) {
-    auto* value_ty = TypeOf(value)->UnwrapRef();
+    auto* value_ty = sem_.TypeOf(value)->UnwrapRef();
     if (value_ty->is_scalar()) {
       if (elem_ty != value_ty) {
         AddError(
             "type in vector constructor does not match vector type: "
             "expected '" +
-                TypeNameOf(elem_ty) + "', found '" + TypeNameOf(value_ty) + "'",
+                sem_.TypeNameOf(elem_ty) + "', found '" +
+                sem_.TypeNameOf(value_ty) + "'",
             value->source);
         return false;
       }
@@ -1738,8 +1755,8 @@
         AddError(
             "type in vector constructor does not match vector type: "
             "expected '" +
-                TypeNameOf(elem_ty) + "', found '" + TypeNameOf(value_elem_ty) +
-                "'",
+                sem_.TypeNameOf(elem_ty) + "', found '" +
+                sem_.TypeNameOf(value_elem_ty) + "'",
             value->source);
         return false;
       }
@@ -1748,7 +1765,7 @@
     } else {
       // A vector constructor can only accept vectors and scalars.
       AddError("expected vector or scalar type in vector constructor; found: " +
-                   TypeNameOf(value_ty),
+                   sem_.TypeNameOf(value_ty),
                value->source);
       return false;
     }
@@ -1764,15 +1781,17 @@
     }
     const Source& values_start = values[0]->source;
     const Source& values_end = values[values.size() - 1]->source;
-    AddError("attempted to construct '" + TypeNameOf(vec_type) + "' with " +
-                 std::to_string(value_cardinality_sum) + " component(s)",
+    AddError("attempted to construct '" + sem_.TypeNameOf(vec_type) +
+                 "' with " + std::to_string(value_cardinality_sum) +
+                 " component(s)",
              Source::Combine(values_start, values_end));
     return false;
   }
   return true;
 }
 
-bool Resolver::ValidateVector(const sem::Vector* ty, const Source& source) {
+bool Resolver::ValidateVector(const sem::Vector* ty,
+                              const Source& source) const {
   if (!ty->type()->is_scalar()) {
     AddError("vector element type must be 'bool', 'f32', 'i32' or 'u32'",
              source);
@@ -1781,7 +1800,8 @@
   return true;
 }
 
-bool Resolver::ValidateMatrix(const sem::Matrix* ty, const Source& source) {
+bool Resolver::ValidateMatrix(const sem::Matrix* ty,
+                              const Source& source) const {
   if (!ty->is_float_matrix()) {
     AddError("matrix element type must be 'f32'", source);
     return false;
@@ -1789,8 +1809,9 @@
   return true;
 }
 
-bool Resolver::ValidateMatrixConstructorOrCast(const ast::CallExpression* ctor,
-                                               const sem::Matrix* matrix_ty) {
+bool Resolver::ValidateMatrixConstructorOrCast(
+    const ast::CallExpression* ctor,
+    const sem::Matrix* matrix_ty) const {
   auto& values = ctor->args;
   // Zero Value expression
   if (values.empty()) {
@@ -1804,7 +1825,7 @@
   std::vector<const sem::Type*> arg_tys;
   arg_tys.reserve(values.size());
   for (auto* value : values) {
-    arg_tys.emplace_back(TypeOf(value)->UnwrapRef());
+    arg_tys.emplace_back(sem_.TypeOf(value)->UnwrapRef());
   }
 
   auto* elem_type = matrix_ty->type();
@@ -1815,8 +1836,8 @@
   auto print_error = [&]() {
     const Source& values_start = values[0]->source;
     const Source& values_end = values[values.size() - 1]->source;
-    auto type_name = TypeNameOf(matrix_ty);
-    auto elem_type_name = TypeNameOf(elem_type);
+    auto type_name = sem_.TypeNameOf(matrix_ty);
+    auto elem_type_name = sem_.TypeNameOf(elem_type);
     std::stringstream ss;
     ss << "no matching constructor " + type_name << "(";
     for (size_t i = 0; i < values.size(); i++) {
@@ -1865,7 +1886,7 @@
 }
 
 bool Resolver::ValidateScalarConstructorOrCast(const ast::CallExpression* ctor,
-                                               const sem::Type* ty) {
+                                               const sem::Type* ty) const {
   if (ctor->args.size() == 0) {
     return true;
   }
@@ -1878,7 +1899,7 @@
 
   // Validate constructor
   auto* value = ctor->args[0];
-  auto* value_ty = TypeOf(value)->UnwrapRef();
+  auto* value_ty = sem_.TypeOf(value)->UnwrapRef();
 
   using Bool = sem::Bool;
   using I32 = sem::I32;
@@ -1890,8 +1911,8 @@
                         (ty->Is<U32>() && value_ty->is_scalar()) ||
                         (ty->Is<F32>() && value_ty->is_scalar());
   if (!is_valid) {
-    AddError("cannot construct '" + TypeNameOf(ty) +
-                 "' with a value of type '" + TypeNameOf(value_ty) + "'",
+    AddError("cannot construct '" + sem_.TypeNameOf(ty) +
+                 "' with a value of type '" + sem_.TypeNameOf(value_ty) + "'",
              ctor->source);
 
     return false;
@@ -1900,7 +1921,7 @@
   return true;
 }
 
-bool Resolver::ValidatePipelineStages() {
+bool Resolver::ValidatePipelineStages() const {
   auto check_workgroup_storage = [&](const sem::Function* func,
                                      const sem::Function* entry_point) {
     auto stage = entry_point->Declaration()->PipelineStage();
@@ -1995,7 +2016,8 @@
   return true;
 }
 
-bool Resolver::ValidateArray(const sem::Array* arr, const Source& source) {
+bool Resolver::ValidateArray(const sem::Array* arr,
+                             const Source& source) const {
   auto* el_ty = arr->ElemType();
 
   if (!IsFixedFootprint(el_ty)) {
@@ -2009,7 +2031,7 @@
 bool Resolver::ValidateArrayStrideAttribute(const ast::StrideAttribute* attr,
                                             uint32_t el_size,
                                             uint32_t el_align,
-                                            const Source& source) {
+                                            const Source& source) const {
   auto stride = attr->stride;
   bool is_valid_stride =
       (stride >= el_size) && (stride >= el_align) && (stride % el_align == 0);
@@ -2028,7 +2050,7 @@
   return true;
 }
 
-bool Resolver::ValidateAlias(const ast::Alias* alias) {
+bool Resolver::ValidateAlias(const ast::Alias* alias) const {
   auto name = builder_->Symbols().NameFor(alias->name);
   if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
     AddError("'" + name + "' is a builtin and cannot be redeclared as an alias",
@@ -2039,7 +2061,8 @@
   return true;
 }
 
-bool Resolver::ValidateStructure(const sem::Struct* str) {
+bool Resolver::ValidateStructure(const sem::Struct* str,
+                                 ast::PipelineStage stage) const {
   auto name = builder_->Symbols().NameFor(str->Declaration()->name);
   if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
     AddError("'" + name + "' is a builtin and cannot be redeclared as a struct",
@@ -2100,11 +2123,11 @@
       } else if (auto* location = attr->As<ast::LocationAttribute>()) {
         has_location = true;
         if (!ValidateLocationAttribute(location, member->Type(), locations,
-                                       member->Declaration()->source)) {
+                                       stage, member->Declaration()->source)) {
           return false;
         }
       } else if (auto* builtin = attr->As<ast::BuiltinAttribute>()) {
-        if (!ValidateBuiltinAttribute(builtin, member->Type(),
+        if (!ValidateBuiltinAttribute(builtin, member->Type(), stage,
                                       /* is_input */ false)) {
           return false;
         }
@@ -2146,18 +2169,18 @@
     const ast::LocationAttribute* location,
     const sem::Type* type,
     std::unordered_set<uint32_t>& locations,
+    ast::PipelineStage stage,
     const Source& source,
-    const bool is_input) {
+    const bool is_input) const {
   std::string inputs_or_output = is_input ? "inputs" : "output";
-  if (current_function_ && current_function_->Declaration()->PipelineStage() ==
-                               ast::PipelineStage::kCompute) {
+  if (stage == ast::PipelineStage::kCompute) {
     AddError("attribute is not valid for compute shader " + inputs_or_output,
              location->source);
     return false;
   }
 
   if (!type->is_numeric_scalar_or_vector()) {
-    std::string invalid_type = TypeNameOf(type);
+    std::string invalid_type = sem_.TypeNameOf(type);
     AddError("cannot apply 'location' attribute to declaration of type '" +
                  invalid_type + "'",
              source);
@@ -2178,23 +2201,20 @@
   return true;
 }
 
-bool Resolver::ValidateReturn(const ast::ReturnStatement* ret) {
-  auto* func_type = current_function_->ReturnType();
-
-  auto* ret_type = ret->value ? TypeOf(ret->value)->UnwrapRef()
-                              : builder_->create<sem::Void>();
-
+bool Resolver::ValidateReturn(const ast::ReturnStatement* ret,
+                              const sem::Type* func_type,
+                              const sem::Type* ret_type) const {
   if (func_type->UnwrapRef() != ret_type) {
     AddError(
         "return statement type must match its function "
         "return type, returned '" +
-            TypeNameOf(ret_type) + "', expected '" + TypeNameOf(func_type) +
-            "'",
+            sem_.TypeNameOf(ret_type) + "', expected '" +
+            sem_.TypeNameOf(func_type) + "'",
         ret->source);
     return false;
   }
 
-  auto* sem = Sem(ret);
+  auto* sem = sem_.Get(ret);
   if (auto* continuing = ClosestContinuing(/*stop_at_loop*/ false)) {
     AddError("continuing blocks must not contain a return statement",
              ret->source);
@@ -2209,7 +2229,7 @@
 }
 
 bool Resolver::ValidateSwitch(const ast::SwitchStatement* s) {
-  auto* cond_ty = TypeOf(s->condition)->UnwrapRef();
+  auto* cond_ty = sem_.TypeOf(s->condition)->UnwrapRef();
   if (!cond_ty->is_integer_scalar()) {
     AddError(
         "switch statement selector expression must be of a "
@@ -2233,7 +2253,7 @@
     }
 
     for (auto* selector : case_stmt->selectors) {
-      if (cond_ty != TypeOf(selector)) {
+      if (cond_ty != sem_.TypeOf(selector)) {
         AddError(
             "the case selector values must have the same "
             "type as the selector expression.",
@@ -2265,7 +2285,7 @@
 }
 
 bool Resolver::ValidateAssignment(const ast::Statement* a,
-                                  const sem::Type* rhs_ty) {
+                                  const sem::Type* rhs_ty) const {
   const ast::Expression* lhs;
   const ast::Expression* rhs;
   if (auto* assign = a->As<ast::AssignmentStatement>()) {
@@ -2285,7 +2305,7 @@
     if (!ty->IsConstructible() &&
         !ty->IsAnyOf<sem::Pointer, sem::Texture, sem::Sampler>()) {
       AddError(
-          "cannot assign '" + TypeNameOf(rhs_ty) +
+          "cannot assign '" + sem_.TypeNameOf(rhs_ty) +
               "' to '_'. '_' can only be assigned a constructible, pointer, "
               "texture or sampler type",
           rhs->source);
@@ -2295,7 +2315,7 @@
   }
 
   // https://gpuweb.github.io/gpuweb/wgsl/#assignment-statement
-  auto const* lhs_ty = TypeOf(lhs);
+  auto const* lhs_ty = sem_.TypeOf(lhs);
 
   if (auto* var = ResolvedSymbol<sem::Variable>(lhs)) {
     auto* decl = var->Declaration();
@@ -2318,7 +2338,7 @@
   auto* lhs_ref = lhs_ty->As<sem::Reference>();
   if (!lhs_ref) {
     // LHS is not a reference, so it has no storage.
-    AddError("cannot assign to value of type '" + TypeNameOf(lhs_ty) + "'",
+    AddError("cannot assign to value of type '" + sem_.TypeNameOf(lhs_ty) + "'",
              lhs->source);
     return false;
   }
@@ -2328,8 +2348,8 @@
 
   // Value type has to match storage type
   if (storage_ty != value_type) {
-    AddError("cannot assign '" + TypeNameOf(rhs_ty) + "' to '" +
-                 TypeNameOf(lhs_ty) + "'",
+    AddError("cannot assign '" + sem_.TypeNameOf(rhs_ty) + "' to '" +
+                 sem_.TypeNameOf(lhs_ty) + "'",
              a->source);
     return false;
   }
@@ -2338,16 +2358,16 @@
     return false;
   }
   if (lhs_ref->Access() == ast::Access::kRead) {
-    AddError(
-        "cannot store into a read-only type '" + RawTypeNameOf(lhs_ty) + "'",
-        a->source);
+    AddError("cannot store into a read-only type '" +
+                 sem_.RawTypeNameOf(lhs_ty) + "'",
+             a->source);
     return false;
   }
   return true;
 }
 
 bool Resolver::ValidateIncrementDecrementStatement(
-    const ast::IncrementDecrementStatement* inc) {
+    const ast::IncrementDecrementStatement* inc) const {
   const ast::Expression* lhs = inc->lhs;
 
   // https://gpuweb.github.io/gpuweb/wgsl/#increment-decrement
@@ -2370,11 +2390,11 @@
     }
   }
 
-  auto const* lhs_ty = TypeOf(lhs);
+  auto const* lhs_ty = sem_.TypeOf(lhs);
   auto* lhs_ref = lhs_ty->As<sem::Reference>();
   if (!lhs_ref) {
     // LHS is not a reference, so it has no storage.
-    AddError("cannot modify value of type '" + TypeNameOf(lhs_ty) + "'",
+    AddError("cannot modify value of type '" + sem_.TypeNameOf(lhs_ty) + "'",
              lhs->source);
     return false;
   }
@@ -2387,15 +2407,16 @@
   }
 
   if (lhs_ref->Access() == ast::Access::kRead) {
-    AddError("cannot modify read-only type '" + RawTypeNameOf(lhs_ty) + "'",
-             inc->source);
+    AddError(
+        "cannot modify read-only type '" + sem_.RawTypeNameOf(lhs_ty) + "'",
+        inc->source);
     return false;
   }
   return true;
 }
 
 bool Resolver::ValidateNoDuplicateAttributes(
-    const ast::AttributeList& attributes) {
+    const ast::AttributeList& attributes) const {
   std::unordered_map<const TypeInfo*, Source> seen;
   for (auto* d : attributes) {
     auto res = seen.emplace(&d->TypeInfo(), d->source);
@@ -2425,4 +2446,10 @@
   return !IsValidationDisabled(attributes, validation);
 }
 
+std::string Resolver::VectorPretty(uint32_t size,
+                                   const sem::Type* element_type) const {
+  sem::Vector vec_type(element_type, size);
+  return vec_type.FriendlyName(builder_->Symbols());
+}
+
 }  // namespace tint::resolver
diff --git a/src/tint/resolver/sem_helper.cc b/src/tint/resolver/sem_helper.cc
new file mode 100644
index 0000000..74b3c5b
--- /dev/null
+++ b/src/tint/resolver/sem_helper.cc
@@ -0,0 +1,60 @@
+// Copyright 2022 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/tint/resolver/sem_helper.h"
+
+#include "src/tint/sem/expression.h"
+
+namespace tint::resolver {
+
+SemHelper::SemHelper(ProgramBuilder* builder) : builder_(builder) {}
+
+SemHelper::~SemHelper() = default;
+
+std::string SemHelper::TypeNameOf(const sem::Type* ty) const {
+  return RawTypeNameOf(ty->UnwrapRef());
+}
+
+std::string SemHelper::RawTypeNameOf(const sem::Type* ty) const {
+  return ty->FriendlyName(builder_->Symbols());
+}
+
+sem::Type* SemHelper::TypeOf(const ast::Expression* expr) const {
+  auto* sem = Get(expr);
+  return sem ? const_cast<sem::Type*>(sem->Type()) : nullptr;
+}
+
+sem::Type* SemHelper::TypeOf(const ast::LiteralExpression* lit) {
+  return Switch(
+      lit,
+      [&](const ast::SintLiteralExpression*) {
+        return builder_->create<sem::I32>();
+      },
+      [&](const ast::UintLiteralExpression*) {
+        return builder_->create<sem::U32>();
+      },
+      [&](const ast::FloatLiteralExpression*) {
+        return builder_->create<sem::F32>();
+      },
+      [&](const ast::BoolLiteralExpression*) {
+        return builder_->create<sem::Bool>();
+      },
+      [&](Default) {
+        TINT_UNREACHABLE(Resolver, builder_->Diagnostics())
+            << "Unhandled literal type: " << lit->TypeInfo().name;
+        return nullptr;
+      });
+}
+
+}  // namespace tint::resolver
diff --git a/src/tint/resolver/sem_helper.h b/src/tint/resolver/sem_helper.h
new file mode 100644
index 0000000..0e95397
--- /dev/null
+++ b/src/tint/resolver/sem_helper.h
@@ -0,0 +1,74 @@
+// Copyright 2022 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_TINT_RESOLVER_SEM_HELPER_H_
+#define SRC_TINT_RESOLVER_SEM_HELPER_H_
+
+#include <string>
+
+#include "src/tint/diagnostic/diagnostic.h"
+#include "src/tint/program_builder.h"
+
+namespace tint::resolver {
+
+/// Helper class to retrieve sem information.
+class SemHelper {
+ public:
+  /// Constructor
+  /// @param builder the program builder
+  explicit SemHelper(ProgramBuilder* builder);
+  ~SemHelper();
+
+  /// Get is a helper for obtaining the semantic node for the given AST node.
+  /// @param ast the ast node to get the sem for
+  /// @returns the sem node for the provided |ast|
+  template <typename SEM = sem::Info::InferFromAST,
+            typename AST_OR_TYPE = CastableBase>
+  auto* Get(const AST_OR_TYPE* ast) const {
+    using T = sem::Info::GetResultType<SEM, AST_OR_TYPE>;
+    auto* sem = builder_->Sem().Get(ast);
+    if (!sem) {
+      TINT_ICE(Resolver, builder_->Diagnostics())
+          << "AST node '" << ast->TypeInfo().name << "' had no semantic info\n"
+          << "At: " << ast->source << "\n"
+          << "Pointer: " << ast;
+    }
+    return const_cast<T*>(As<T>(sem));
+  }
+
+  /// @returns the resolved type of the ast::Expression `expr`
+  /// @param expr the expression
+  sem::Type* TypeOf(const ast::Expression* expr) const;
+
+  /// @returns the semantic type of the AST literal `lit`
+  /// @param lit the literal
+  sem::Type* TypeOf(const ast::LiteralExpression* lit);
+
+  /// @returns the type name of the given semantic type, unwrapping
+  /// references.
+  /// @param ty the type to look up
+  std::string TypeNameOf(const sem::Type* ty) const;
+
+  /// @returns the type name of the given semantic type, without unwrapping
+  /// references.
+  /// @param ty the type to look up
+  std::string RawTypeNameOf(const sem::Type* ty) const;
+
+ private:
+  ProgramBuilder* builder_;
+};
+
+}  // namespace tint::resolver
+
+#endif  // SRC_TINT_RESOLVER_SEM_HELPER_H_