diff --git a/fuzzers/tint_common_fuzzer.cc b/fuzzers/tint_common_fuzzer.cc
index 4334fc7..6c92306 100644
--- a/fuzzers/tint_common_fuzzer.cc
+++ b/fuzzers/tint_common_fuzzer.cc
@@ -35,12 +35,13 @@
   __builtin_trap();
 }
 
-[[noreturn]] void ValidityErrorReporter() {
+[[noreturn]] void ValidityErrorReporter(const tint::diag::List& diags) {
   auto printer = tint::diag::Printer::create(stderr, true);
   printer->write(
       "Fuzzing detected valid input program being transformed into an invalid "
       "output progam\n",
       {diag::Color::kRed, true});
+  tint::diag::Formatter().format(diags, printer.get());
   __builtin_trap();
 }
 
@@ -258,7 +259,14 @@
   if (transform_manager_) {
     auto out = transform_manager_->Run(&program, transform_inputs_);
     if (!out.program.IsValid()) {
-      ValidityErrorReporter();
+      // Transforms can produce error messages for bad input.
+      // Catch ICEs and errors from non transform systems.
+      for (auto diag : out.program.Diagnostics()) {
+        if (diag.severity > diag::Severity::Error ||
+            diag.system != diag::System::Transform) {
+          ValidityErrorReporter(out.program.Diagnostics());
+        }
+      }
     }
 
     program = std::move(out.program);
diff --git a/src/ast/alias.cc b/src/ast/alias.cc
index be02bb2..ffa32ac 100644
--- a/src/ast/alias.cc
+++ b/src/ast/alias.cc
@@ -28,7 +28,7 @@
     : Base(program_id, source, name),
       subtype_(subtype),
       type_name_("__alias_" + name.to_str() + subtype->type_name()) {
-  TINT_ASSERT(subtype_);
+  TINT_ASSERT(AST, subtype_);
 }
 
 Alias::Alias(Alias&&) = default;
diff --git a/src/ast/array.cc b/src/ast/array.cc
index 96ac1de..2da027f 100644
--- a/src/ast/array.cc
+++ b/src/ast/array.cc
@@ -38,7 +38,7 @@
 Array::~Array() = default;
 
 std::string Array::type_name() const {
-  TINT_ASSERT(subtype_);
+  TINT_ASSERT(AST, subtype_);
 
   std::string type_name = "__array" + subtype_->type_name();
   if (!IsRuntimeArray()) {
diff --git a/src/ast/array_accessor_expression.cc b/src/ast/array_accessor_expression.cc
index bb358b6..23fc1b2 100644
--- a/src/ast/array_accessor_expression.cc
+++ b/src/ast/array_accessor_expression.cc
@@ -26,10 +26,10 @@
                                                  Expression* array,
                                                  Expression* idx_expr)
     : Base(program_id, source), array_(array), idx_expr_(idx_expr) {
-  TINT_ASSERT(array_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(array_, program_id);
-  TINT_ASSERT(idx_expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(idx_expr_, program_id);
+  TINT_ASSERT(AST, array_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, array_, program_id);
+  TINT_ASSERT(AST, idx_expr_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, idx_expr_, program_id);
 }
 
 ArrayAccessorExpression::ArrayAccessorExpression(ArrayAccessorExpression&&) =
diff --git a/src/ast/assignment_statement.cc b/src/ast/assignment_statement.cc
index 9f672a2..7af726b 100644
--- a/src/ast/assignment_statement.cc
+++ b/src/ast/assignment_statement.cc
@@ -26,10 +26,10 @@
                                          Expression* lhs,
                                          Expression* rhs)
     : Base(program_id, source), lhs_(lhs), rhs_(rhs) {
-  TINT_ASSERT(lhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(lhs_, program_id);
-  TINT_ASSERT(rhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(rhs_, program_id);
+  TINT_ASSERT(AST, lhs_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs_, program_id);
+  TINT_ASSERT(AST, rhs_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs_, program_id);
 }
 
 AssignmentStatement::AssignmentStatement(AssignmentStatement&&) = default;
diff --git a/src/ast/binary_expression.cc b/src/ast/binary_expression.cc
index 8a6a1c2..6da0c57 100644
--- a/src/ast/binary_expression.cc
+++ b/src/ast/binary_expression.cc
@@ -27,11 +27,11 @@
                                    Expression* lhs,
                                    Expression* rhs)
     : Base(program_id, source), op_(op), lhs_(lhs), rhs_(rhs) {
-  TINT_ASSERT(lhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(lhs_, program_id);
-  TINT_ASSERT(rhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(rhs_, program_id);
-  TINT_ASSERT(op_ != BinaryOp::kNone);
+  TINT_ASSERT(AST, lhs_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs_, program_id);
+  TINT_ASSERT(AST, rhs_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs_, program_id);
+  TINT_ASSERT(AST, op_ != BinaryOp::kNone);
 }
 
 BinaryExpression::BinaryExpression(BinaryExpression&&) = default;
diff --git a/src/ast/bitcast_expression.cc b/src/ast/bitcast_expression.cc
index 369d6ae..aafa160 100644
--- a/src/ast/bitcast_expression.cc
+++ b/src/ast/bitcast_expression.cc
@@ -26,9 +26,9 @@
                                      ast::Type* type,
                                      Expression* expr)
     : Base(program_id, source), type_(type), expr_(expr) {
-  TINT_ASSERT(type_);
-  TINT_ASSERT(expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(expr, program_id);
+  TINT_ASSERT(AST, type_);
+  TINT_ASSERT(AST, expr_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, expr, program_id);
 }
 
 BitcastExpression::BitcastExpression(BitcastExpression&&) = default;
diff --git a/src/ast/block_statement.cc b/src/ast/block_statement.cc
index 9a4cbf0..c1d6e0b 100644
--- a/src/ast/block_statement.cc
+++ b/src/ast/block_statement.cc
@@ -26,8 +26,8 @@
                                const StatementList& statements)
     : Base(program_id, source), statements_(std::move(statements)) {
   for (auto* stmt : *this) {
-    TINT_ASSERT(stmt);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(stmt, program_id);
+    TINT_ASSERT(AST, stmt);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, stmt, program_id);
   }
 }
 
diff --git a/src/ast/call_expression.cc b/src/ast/call_expression.cc
index e63b41f..752af8f 100644
--- a/src/ast/call_expression.cc
+++ b/src/ast/call_expression.cc
@@ -26,11 +26,11 @@
                                IdentifierExpression* func,
                                ExpressionList params)
     : Base(program_id, source), func_(func), params_(params) {
-  TINT_ASSERT(func_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(func_, program_id);
+  TINT_ASSERT(AST, func_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, func_, program_id);
   for (auto* param : params_) {
-    TINT_ASSERT(param);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(param, program_id);
+    TINT_ASSERT(AST, param);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, param, program_id);
   }
 }
 
diff --git a/src/ast/call_statement.cc b/src/ast/call_statement.cc
index c771a81..83361f6 100644
--- a/src/ast/call_statement.cc
+++ b/src/ast/call_statement.cc
@@ -25,8 +25,8 @@
                              const Source& source,
                              CallExpression* call)
     : Base(program_id, source), call_(call) {
-  TINT_ASSERT(call_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(call_, program_id);
+  TINT_ASSERT(AST, call_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, call_, program_id);
 }
 
 CallStatement::CallStatement(CallStatement&&) = default;
diff --git a/src/ast/case_statement.cc b/src/ast/case_statement.cc
index 90c8309..0002e03 100644
--- a/src/ast/case_statement.cc
+++ b/src/ast/case_statement.cc
@@ -26,11 +26,11 @@
                              CaseSelectorList selectors,
                              BlockStatement* body)
     : Base(program_id, source), selectors_(selectors), body_(body) {
-  TINT_ASSERT(body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT(AST, body_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
   for (auto* selector : selectors) {
-    TINT_ASSERT(selector);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(selector, program_id);
+    TINT_ASSERT(AST, selector);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, selector, program_id);
   }
 }
 
diff --git a/src/ast/depth_texture.cc b/src/ast/depth_texture.cc
index b669113..f33f9aa 100644
--- a/src/ast/depth_texture.cc
+++ b/src/ast/depth_texture.cc
@@ -33,7 +33,7 @@
                            const Source& source,
                            TextureDimension dim)
     : Base(program_id, source, dim) {
-  TINT_ASSERT(IsValidDepthDimension(dim));
+  TINT_ASSERT(AST, IsValidDepthDimension(dim));
 }
 
 DepthTexture::DepthTexture(DepthTexture&&) = default;
diff --git a/src/ast/else_statement.cc b/src/ast/else_statement.cc
index 7e43364..40271d3 100644
--- a/src/ast/else_statement.cc
+++ b/src/ast/else_statement.cc
@@ -26,9 +26,9 @@
                              Expression* condition,
                              BlockStatement* body)
     : Base(program_id, source), condition_(condition), body_(body) {
-  TINT_ASSERT(body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
+  TINT_ASSERT(AST, body_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
 }
 
 ElseStatement::ElseStatement(ElseStatement&&) = default;
diff --git a/src/ast/function.cc b/src/ast/function.cc
index 0896e7a..c0014e1 100644
--- a/src/ast/function.cc
+++ b/src/ast/function.cc
@@ -38,19 +38,19 @@
       body_(body),
       decorations_(std::move(decorations)),
       return_type_decorations_(std::move(return_type_decorations)) {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(symbol_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, symbol_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body, program_id);
   for (auto* param : params_) {
-    TINT_ASSERT(param && param->is_const());
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(param, program_id);
+    TINT_ASSERT(AST, param && param->is_const());
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, param, program_id);
   }
-  TINT_ASSERT(symbol_.IsValid());
-  TINT_ASSERT(return_type_);
+  TINT_ASSERT(AST, symbol_.IsValid());
+  TINT_ASSERT(AST, return_type_);
   for (auto* deco : decorations_) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
   for (auto* deco : return_type_decorations_) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
 }
 
diff --git a/src/ast/identifier_expression.cc b/src/ast/identifier_expression.cc
index fc32c43..df9e084 100644
--- a/src/ast/identifier_expression.cc
+++ b/src/ast/identifier_expression.cc
@@ -25,8 +25,8 @@
                                            const Source& source,
                                            Symbol sym)
     : Base(program_id, source), sym_(sym) {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(sym_, program_id);
-  TINT_ASSERT(sym_.IsValid());
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, sym_, program_id);
+  TINT_ASSERT(AST, sym_.IsValid());
 }
 
 IdentifierExpression::IdentifierExpression(IdentifierExpression&&) = default;
diff --git a/src/ast/if_statement.cc b/src/ast/if_statement.cc
index 8ddfe73..b0daf49 100644
--- a/src/ast/if_statement.cc
+++ b/src/ast/if_statement.cc
@@ -30,13 +30,13 @@
       condition_(condition),
       body_(body),
       else_statements_(std::move(else_stmts)) {
-  TINT_ASSERT(condition_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
-  TINT_ASSERT(body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT(AST, condition_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
+  TINT_ASSERT(AST, body_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
   for (auto* el : else_statements_) {
-    TINT_ASSERT(el);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(el, program_id);
+    TINT_ASSERT(AST, el);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, el, program_id);
   }
 }
 
diff --git a/src/ast/intrinsic_texture_helper_test.cc b/src/ast/intrinsic_texture_helper_test.cc
index 3ff6360..20c5f4c 100644
--- a/src/ast/intrinsic_texture_helper_test.cc
+++ b/src/ast/intrinsic_texture_helper_test.cc
@@ -142,7 +142,7 @@
       return b->ty.i32();
   }
 
-  TINT_UNREACHABLE(b->Diagnostics());
+  TINT_UNREACHABLE(AST, b->Diagnostics());
   return {};
 }
 
@@ -176,7 +176,7 @@
     }
   }
 
-  TINT_UNREACHABLE(b->Diagnostics());
+  TINT_UNREACHABLE(AST, b->Diagnostics());
   return nullptr;
 }
 
diff --git a/src/ast/loop_statement.cc b/src/ast/loop_statement.cc
index 67d9701..e6e2134 100644
--- a/src/ast/loop_statement.cc
+++ b/src/ast/loop_statement.cc
@@ -26,9 +26,9 @@
                              BlockStatement* body,
                              BlockStatement* continuing)
     : Base(program_id, source), body_(body), continuing_(continuing) {
-  TINT_ASSERT(body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(continuing_, program_id);
+  TINT_ASSERT(AST, body_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, continuing_, program_id);
 }
 
 LoopStatement::LoopStatement(LoopStatement&&) = default;
diff --git a/src/ast/matrix.cc b/src/ast/matrix.cc
index da8e019..3564513 100644
--- a/src/ast/matrix.cc
+++ b/src/ast/matrix.cc
@@ -30,10 +30,10 @@
       subtype_(subtype),
       rows_(rows),
       columns_(columns) {
-  TINT_ASSERT(rows > 1);
-  TINT_ASSERT(rows < 5);
-  TINT_ASSERT(columns > 1);
-  TINT_ASSERT(columns < 5);
+  TINT_ASSERT(AST, rows > 1);
+  TINT_ASSERT(AST, rows < 5);
+  TINT_ASSERT(AST, columns > 1);
+  TINT_ASSERT(AST, columns < 5);
 }
 
 Matrix::Matrix(Matrix&&) = default;
diff --git a/src/ast/member_accessor_expression.cc b/src/ast/member_accessor_expression.cc
index 1b5e99e..7a25750 100644
--- a/src/ast/member_accessor_expression.cc
+++ b/src/ast/member_accessor_expression.cc
@@ -26,10 +26,10 @@
                                                    Expression* structure,
                                                    IdentifierExpression* member)
     : Base(program_id, source), struct_(structure), member_(member) {
-  TINT_ASSERT(struct_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(struct_, program_id);
-  TINT_ASSERT(member_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(member_, program_id);
+  TINT_ASSERT(AST, struct_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, struct_, program_id);
+  TINT_ASSERT(AST, member_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, member_, program_id);
 }
 
 MemberAccessorExpression::MemberAccessorExpression(MemberAccessorExpression&&) =
diff --git a/src/ast/module.cc b/src/ast/module.cc
index 4aeb996..894d435 100644
--- a/src/ast/module.cc
+++ b/src/ast/module.cc
@@ -44,7 +44,7 @@
       global_variables_.push_back(var);
     } else {
       diag::List diagnostics;
-      TINT_ICE(diagnostics) << "Unknown global declaration type";
+      TINT_ICE(AST, diagnostics) << "Unknown global declaration type";
     }
   }
 }
@@ -61,22 +61,22 @@
 }
 
 void Module::AddGlobalVariable(ast::Variable* var) {
-  TINT_ASSERT(var);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(var, program_id());
+  TINT_ASSERT(AST, var);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, var, program_id());
   global_variables_.push_back(var);
   global_declarations_.push_back(var);
 }
 
 void Module::AddTypeDecl(ast::TypeDecl* type) {
-  TINT_ASSERT(type);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(type, program_id());
+  TINT_ASSERT(AST, type);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, type, program_id());
   type_decls_.push_back(type);
   global_declarations_.push_back(type);
 }
 
 void Module::AddFunction(ast::Function* func) {
-  TINT_ASSERT(func);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(func, program_id());
+  TINT_ASSERT(AST, func);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, func, program_id());
   functions_.push_back(func);
   global_declarations_.push_back(func);
 }
@@ -98,20 +98,22 @@
 
   for (auto* decl : global_declarations_) {
     if (!decl) {
-      TINT_ICE(ctx->dst->Diagnostics()) << "src global declaration was nullptr";
+      TINT_ICE(AST, ctx->dst->Diagnostics())
+          << "src global declaration was nullptr";
       continue;
     }
     if (auto* type = decl->As<ast::TypeDecl>()) {
-      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(type, program_id());
+      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, type, program_id());
       type_decls_.push_back(type);
     } else if (auto* func = decl->As<Function>()) {
-      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(func, program_id());
+      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, func, program_id());
       functions_.push_back(func);
     } else if (auto* var = decl->As<Variable>()) {
-      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(var, program_id());
+      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, var, program_id());
       global_variables_.push_back(var);
     } else {
-      TINT_ICE(ctx->dst->Diagnostics()) << "Unknown global declaration type";
+      TINT_ICE(AST, ctx->dst->Diagnostics())
+          << "Unknown global declaration type";
     }
   }
 }
diff --git a/src/ast/multisampled_texture.cc b/src/ast/multisampled_texture.cc
index 5f907ab..51048d3 100644
--- a/src/ast/multisampled_texture.cc
+++ b/src/ast/multisampled_texture.cc
@@ -26,7 +26,7 @@
                                          TextureDimension dim,
                                          Type* type)
     : Base(program_id, source, dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
 }
 
 MultisampledTexture::MultisampledTexture(MultisampledTexture&&) = default;
diff --git a/src/ast/return_statement.cc b/src/ast/return_statement.cc
index e545136..a8e1bfe 100644
--- a/src/ast/return_statement.cc
+++ b/src/ast/return_statement.cc
@@ -28,7 +28,7 @@
                                  const Source& source,
                                  Expression* value)
     : Base(program_id, source), value_(value) {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(value_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, value_, program_id);
 }
 
 ReturnStatement::ReturnStatement(ReturnStatement&&) = default;
diff --git a/src/ast/sampled_texture.cc b/src/ast/sampled_texture.cc
index 6fe7cff..69cad40 100644
--- a/src/ast/sampled_texture.cc
+++ b/src/ast/sampled_texture.cc
@@ -26,7 +26,7 @@
                                TextureDimension dim,
                                Type const* type)
     : Base(program_id, source, dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
 }
 
 SampledTexture::SampledTexture(SampledTexture&&) = default;
diff --git a/src/ast/scalar_constructor_expression.cc b/src/ast/scalar_constructor_expression.cc
index c8bba78..d74e681 100644
--- a/src/ast/scalar_constructor_expression.cc
+++ b/src/ast/scalar_constructor_expression.cc
@@ -25,8 +25,8 @@
                                                          const Source& source,
                                                          Literal* literal)
     : Base(program_id, source), literal_(literal) {
-  TINT_ASSERT(literal);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(literal, program_id);
+  TINT_ASSERT(AST, literal);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, literal, program_id);
 }
 
 ScalarConstructorExpression::ScalarConstructorExpression(
diff --git a/src/ast/struct.cc b/src/ast/struct.cc
index a59c3a3..9138ee4 100644
--- a/src/ast/struct.cc
+++ b/src/ast/struct.cc
@@ -33,12 +33,12 @@
       members_(std::move(members)),
       decorations_(std::move(decorations)) {
   for (auto* mem : members_) {
-    TINT_ASSERT(mem);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(mem, program_id);
+    TINT_ASSERT(AST, mem);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, mem, program_id);
   }
   for (auto* deco : decorations_) {
-    TINT_ASSERT(deco);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT(AST, deco);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
 }
 
diff --git a/src/ast/struct_member.cc b/src/ast/struct_member.cc
index 11fad7d..743cae0 100644
--- a/src/ast/struct_member.cc
+++ b/src/ast/struct_member.cc
@@ -30,12 +30,12 @@
       symbol_(sym),
       type_(type),
       decorations_(std::move(decorations)) {
-  TINT_ASSERT(type);
-  TINT_ASSERT(symbol_.IsValid());
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(symbol_, program_id);
+  TINT_ASSERT(AST, type);
+  TINT_ASSERT(AST, symbol_.IsValid());
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, symbol_, program_id);
   for (auto* deco : decorations_) {
-    TINT_ASSERT(deco);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT(AST, deco);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
 }
 
diff --git a/src/ast/switch_statement.cc b/src/ast/switch_statement.cc
index b9511f7..2dab086 100644
--- a/src/ast/switch_statement.cc
+++ b/src/ast/switch_statement.cc
@@ -26,11 +26,11 @@
                                  Expression* condition,
                                  CaseStatementList body)
     : Base(program_id, source), condition_(condition), body_(body) {
-  TINT_ASSERT(condition_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
+  TINT_ASSERT(AST, condition_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
   for (auto* stmt : body_) {
-    TINT_ASSERT(stmt);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(stmt, program_id);
+    TINT_ASSERT(AST, stmt);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, stmt, program_id);
   }
 }
 
diff --git a/src/ast/type_constructor_expression.cc b/src/ast/type_constructor_expression.cc
index d0877bf..84e7be4 100644
--- a/src/ast/type_constructor_expression.cc
+++ b/src/ast/type_constructor_expression.cc
@@ -26,10 +26,10 @@
                                                      ast::Type* type,
                                                      ExpressionList values)
     : Base(program_id, source), type_(type), values_(std::move(values)) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
   for (auto* val : values_) {
-    TINT_ASSERT(val);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(val, program_id);
+    TINT_ASSERT(AST, val);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, val, program_id);
   }
 }
 
diff --git a/src/ast/type_decl.cc b/src/ast/type_decl.cc
index 4337d39..566fceb 100644
--- a/src/ast/type_decl.cc
+++ b/src/ast/type_decl.cc
@@ -23,7 +23,7 @@
 
 TypeDecl::TypeDecl(ProgramID program_id, const Source& source, Symbol name)
     : Base(program_id, source), name_(name) {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(name, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, name, program_id);
 }
 
 TypeDecl::TypeDecl(TypeDecl&&) = default;
diff --git a/src/ast/unary_op_expression.cc b/src/ast/unary_op_expression.cc
index be3d5ef..41845a2 100644
--- a/src/ast/unary_op_expression.cc
+++ b/src/ast/unary_op_expression.cc
@@ -26,8 +26,8 @@
                                      UnaryOp op,
                                      Expression* expr)
     : Base(program_id, source), op_(op), expr_(expr) {
-  TINT_ASSERT(expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(expr_, program_id);
+  TINT_ASSERT(AST, expr_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, expr_, program_id);
 }
 
 UnaryOpExpression::UnaryOpExpression(UnaryOpExpression&&) = default;
diff --git a/src/ast/variable.cc b/src/ast/variable.cc
index 6e27a83..be91229 100644
--- a/src/ast/variable.cc
+++ b/src/ast/variable.cc
@@ -40,9 +40,9 @@
       decorations_(std::move(decorations)),
       declared_storage_class_(declared_storage_class),
       declared_access_(declared_access) {
-  TINT_ASSERT(symbol_.IsValid());
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(symbol_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(constructor, program_id);
+  TINT_ASSERT(AST, symbol_.IsValid());
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, symbol_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, constructor, program_id);
 }
 
 Variable::Variable(Variable&&) = default;
diff --git a/src/ast/variable_decl_statement.cc b/src/ast/variable_decl_statement.cc
index 0e41626..549ffea 100644
--- a/src/ast/variable_decl_statement.cc
+++ b/src/ast/variable_decl_statement.cc
@@ -25,8 +25,8 @@
                                              const Source& source,
                                              Variable* variable)
     : Base(program_id, source), variable_(variable) {
-  TINT_ASSERT(variable_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(variable_, program_id);
+  TINT_ASSERT(AST, variable_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, variable_, program_id);
 }
 
 VariableDeclStatement::VariableDeclStatement(VariableDeclStatement&&) = default;
diff --git a/src/ast/vector.cc b/src/ast/vector.cc
index 1b2a8a0..781ef95 100644
--- a/src/ast/vector.cc
+++ b/src/ast/vector.cc
@@ -26,8 +26,8 @@
                Type const* subtype,
                uint32_t size)
     : Base(program_id, source), subtype_(subtype), size_(size) {
-  TINT_ASSERT(size_ > 1);
-  TINT_ASSERT(size_ < 5);
+  TINT_ASSERT(AST, size_ > 1);
+  TINT_ASSERT(AST, size_ < 5);
 }
 
 Vector::Vector(Vector&&) = default;
diff --git a/src/clone_context.h b/src/clone_context.h
index 78b47f6..ff34dd1 100644
--- a/src/clone_context.h
+++ b/src/clone_context.h
@@ -105,7 +105,7 @@
     }
 
     if (src) {
-      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, a);
+      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, a);
     }
 
     // Was Replace() called for this object?
@@ -134,7 +134,7 @@
 
     auto* out = CheckedCast<T>(cloned);
 
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, out);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, out);
 
     return out;
   }
@@ -160,7 +160,7 @@
     }
 
     if (src) {
-      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, a);
+      TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, a);
     }
 
     // Have we seen this object before? If so, return the previously cloned
@@ -323,7 +323,7 @@
     for (auto& transform : transforms_) {
       if (transform.typeinfo->Is(TypeInfo::Of<T>()) ||
           TypeInfo::Of<T>().Is(*transform.typeinfo)) {
-        TINT_ICE(Diagnostics())
+        TINT_ICE(Clone, Diagnostics())
             << "ReplaceAll() called with a handler for type "
             << TypeInfo::Of<T>().name
             << " that is already handled by a handler for type "
@@ -349,8 +349,9 @@
   /// @returns this CloneContext so calls can be chained
   CloneContext& ReplaceAll(const SymbolTransform& replacer) {
     if (symbol_transform_) {
-      TINT_ICE(Diagnostics()) << "ReplaceAll(const SymbolTransform&) called "
-                                 "multiple times on the same CloneContext";
+      TINT_ICE(Clone, Diagnostics())
+          << "ReplaceAll(const SymbolTransform&) called "
+             "multiple times on the same CloneContext";
       return *this;
     }
     symbol_transform_ = replacer;
@@ -369,8 +370,8 @@
   /// @returns this CloneContext so calls can be chained
   template <typename WHAT, typename WITH>
   CloneContext& Replace(WHAT* what, WITH* with) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, what);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, with);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, what);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, with);
     replacements_[what] = with;
     return *this;
   }
@@ -382,9 +383,9 @@
   /// @returns this CloneContext so calls can be chained
   template <typename T, typename OBJECT>
   CloneContext& Remove(const std::vector<T>& vector, OBJECT* object) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, object);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, object);
     if (std::find(vector.begin(), vector.end(), object) == vector.end()) {
-      TINT_ICE(Diagnostics())
+      TINT_ICE(Clone, Diagnostics())
           << "CloneContext::Remove() vector does not contain object";
       return *this;
     }
@@ -400,7 +401,7 @@
   /// @returns this CloneContext so calls can be chained
   template <typename T, typename OBJECT>
   CloneContext& InsertFront(const std::vector<T>& vector, OBJECT* object) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, object);
     auto& transforms = list_transforms_[&vector];
     auto& list = transforms.insert_front_;
     list.emplace_back(object);
@@ -414,7 +415,7 @@
   /// @returns this CloneContext so calls can be chained
   template <typename T, typename OBJECT>
   CloneContext& InsertBack(const std::vector<T>& vector, OBJECT* object) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, object);
     auto& transforms = list_transforms_[&vector];
     auto& list = transforms.insert_back_;
     list.emplace_back(object);
@@ -431,10 +432,10 @@
   CloneContext& InsertBefore(const std::vector<T>& vector,
                              const BEFORE* before,
                              OBJECT* object) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, before);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, before);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, object);
     if (std::find(vector.begin(), vector.end(), before) == vector.end()) {
-      TINT_ICE(Diagnostics())
+      TINT_ICE(Clone, Diagnostics())
           << "CloneContext::InsertBefore() vector does not contain before";
       return *this;
     }
@@ -455,10 +456,10 @@
   CloneContext& InsertAfter(const std::vector<T>& vector,
                             const AFTER* after,
                             OBJECT* object) {
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, after);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, after);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, object);
     if (std::find(vector.begin(), vector.end(), after) == vector.end()) {
-      TINT_ICE(Diagnostics())
+      TINT_ICE(Clone, Diagnostics())
           << "CloneContext::InsertAfter() vector does not contain after";
       return *this;
     }
@@ -505,7 +506,7 @@
     if (TO* cast = As<TO>(obj)) {
       return cast;
     }
-    TINT_ICE(Diagnostics())
+    TINT_ICE(Clone, Diagnostics())
         << "Cloned object was not of the expected type\n"
         << "got:      " << (obj ? obj->TypeInfo().name : "<null>") << "\n"
         << "expected: " << TypeInfo::Of<TO>().name;
diff --git a/src/clone_context_test.cc b/src/clone_context_test.cc
index bd469da..a9865a2 100644
--- a/src/clone_context_test.cc
+++ b/src/clone_context_test.cc
@@ -783,7 +783,7 @@
         Allocator allocator;
         ctx.Clone(allocator.Create<ProgramNode>(ProgramID::New(), dst.ID()));
       },
-      R"(internal compiler error: TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(src, a))");
+      R"(internal compiler error: TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, src, a))");
 }
 
 TEST_F(CloneContextTest, ProgramIDs_Clone_ObjectNotOwnedByDst) {
@@ -795,7 +795,7 @@
         Allocator allocator;
         ctx.Clone(allocator.Create<ProgramNode>(src.ID(), ProgramID::New()));
       },
-      R"(internal compiler error: TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, out))");
+      R"(internal compiler error: TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, out))");
 }
 
 }  // namespace
diff --git a/src/debug.cc b/src/debug.cc
index 7f20d09..03a6fde 100644
--- a/src/debug.cc
+++ b/src/debug.cc
@@ -27,12 +27,13 @@
 
 InternalCompilerError::InternalCompilerError(const char* file,
                                              size_t line,
+                                             diag::System system,
                                              diag::List& diagnostics)
-    : file_(file), line_(line), diagnostics_(diagnostics) {}
+    : file_(file), line_(line), system_(system), diagnostics_(diagnostics) {}
 
 InternalCompilerError::~InternalCompilerError() {
   Source source{Source::Range{Source::Location{line_}}, file_};
-  diagnostics_.add_ice(msg_.str(), source);
+  diagnostics_.add_ice(system_, msg_.str(), source);
 
   if (ice_reporter) {
     ice_reporter(diagnostics_);
diff --git a/src/debug.h b/src/debug.h
index 8c9c7f3..e9bf62b 100644
--- a/src/debug.h
+++ b/src/debug.h
@@ -44,8 +44,12 @@
   /// Constructor
   /// @param file the file containing the ICE
   /// @param line the line containing the ICE
+  /// @param system the Tint system that has raised the ICE
   /// @param diagnostics the list of diagnostics to append the ICE message to
-  InternalCompilerError(const char* file, size_t line, diag::List& diagnostics);
+  InternalCompilerError(const char* file,
+                        size_t line,
+                        diag::System system,
+                        diag::List& diagnostics);
 
   /// Destructor.
   /// Adds the internal compiler error message to the diagnostics list, and then
@@ -64,6 +68,7 @@
  private:
   char const* const file_;
   size_t const line_;
+  diag::System system_;
   diag::List& diagnostics_;
   std::stringstream msg_;
 };
@@ -76,8 +81,9 @@
 /// set.
 /// The ICE message contains the callsite's file and line.
 /// Use the `<<` operator to append an error message to the ICE.
-#define TINT_ICE(diagnostics) \
-  tint::InternalCompilerError(__FILE__, __LINE__, diagnostics)
+#define TINT_ICE(system, diagnostics)             \
+  tint::InternalCompilerError(__FILE__, __LINE__, \
+                              ::tint::diag::System::system, diagnostics)
 
 /// TINT_UNREACHABLE() is a macro for appending a "TINT_UNREACHABLE"
 /// internal compiler error message to the diagnostics list `diagnostics`, and
@@ -85,8 +91,8 @@
 /// reporter is set.
 /// The ICE message contains the callsite's file and line.
 /// Use the `<<` operator to append an error message to the ICE.
-#define TINT_UNREACHABLE(diagnostics) \
-  TINT_ICE(diagnostics) << "TINT_UNREACHABLE "
+#define TINT_UNREACHABLE(system, diagnostics) \
+  TINT_ICE(system, diagnostics) << "TINT_UNREACHABLE "
 
 /// TINT_UNIMPLEMENTED() is a macro for appending a "TINT_UNIMPLEMENTED"
 /// internal compiler error message to the diagnostics list `diagnostics`, and
@@ -94,8 +100,8 @@
 /// reporter is set.
 /// The ICE message contains the callsite's file and line.
 /// Use the `<<` operator to append an error message to the ICE.
-#define TINT_UNIMPLEMENTED(diagnostics) \
-  TINT_ICE(diagnostics) << "TINT_UNIMPLEMENTED "
+#define TINT_UNIMPLEMENTED(system, diagnostics) \
+  TINT_ICE(system, diagnostics) << "TINT_UNIMPLEMENTED "
 
 /// TINT_ASSERT() is a macro for checking the expression is true, triggering a
 /// TINT_ICE if it is not.
@@ -105,12 +111,12 @@
 /// may silently fail in builds where SetInternalCompilerErrorReporter() is not
 /// called. Only use in places where there's no sensible place to put proper
 /// error handling.
-#define TINT_ASSERT(condition)                                      \
-  do {                                                              \
-    if (!(condition)) {                                             \
-      tint::diag::List diagnostics;                                 \
-      TINT_ICE(diagnostics) << "TINT_ASSERT(" << #condition << ")"; \
-    }                                                               \
+#define TINT_ASSERT(system, condition)                                      \
+  do {                                                                      \
+    if (!(condition)) {                                                     \
+      tint::diag::List diagnostics;                                         \
+      TINT_ICE(system, diagnostics) << "TINT_ASSERT(" << #condition << ")"; \
+    }                                                                       \
   } while (false)
 
 #endif  // SRC_DEBUG_H_
diff --git a/src/debug_test.cc b/src/debug_test.cc
index cf3bf1b..6d17c5c 100644
--- a/src/debug_test.cc
+++ b/src/debug_test.cc
@@ -23,17 +23,18 @@
   EXPECT_FATAL_FAILURE(
       {
         diag::List diagnostics;
-        TINT_UNREACHABLE(diagnostics);
+        TINT_UNREACHABLE(Test, diagnostics);
       },
       "internal compiler error");
 }
 
 TEST(DebugTest, AssertTrue) {
-  TINT_ASSERT(true);
+  TINT_ASSERT(Test, true);
 }
 
 TEST(DebugTest, AssertFalse) {
-  EXPECT_FATAL_FAILURE({ TINT_ASSERT(false); }, "internal compiler error");
+  EXPECT_FATAL_FAILURE({ TINT_ASSERT(Test, false); },
+                       "internal compiler error");
 }
 
 }  // namespace
diff --git a/src/diagnostic/diagnostic.h b/src/diagnostic/diagnostic.h
index 0c6847f..3c541ea 100644
--- a/src/diagnostic/diagnostic.h
+++ b/src/diagnostic/diagnostic.h
@@ -32,6 +32,24 @@
   return static_cast<int>(a) >= static_cast<int>(b);
 }
 
+/// System is an enumerator of Tint systems that can be the originator of a
+/// diagnostic message.
+enum class System {
+  AST,
+  Clone,
+  Inspector,
+  Program,
+  ProgramBuilder,
+  Reader,
+  Resolver,
+  Semantic,
+  Symbol,
+  Test,
+  Transform,
+  Utils,
+  Writer,
+};
+
 /// Diagnostic holds all the information for a single compiler diagnostic
 /// message.
 class Diagnostic {
@@ -42,6 +60,8 @@
   Source source;
   /// message is the text associated with the diagnostic.
   std::string message;
+  /// system is the Tint system that raised the diagnostic.
+  System system;
   /// code is the error code, for example a validation error might have the code
   /// `"v-0001"`.
   const char* code = nullptr;
@@ -98,42 +118,54 @@
   }
 
   /// adds the note message with the given Source to the end of this list.
+  /// @param system the system raising the note message
   /// @param note_msg the note message
   /// @param source the source of the note diagnostic
-  void add_note(const std::string& note_msg, const Source& source) {
-    diag::Diagnostic error{};
-    error.severity = diag::Severity::Note;
-    error.source = source;
-    error.message = note_msg;
-    add(std::move(error));
+  void add_note(System system,
+                const std::string& note_msg,
+                const Source& source) {
+    diag::Diagnostic note{};
+    note.severity = diag::Severity::Note;
+    note.system = system;
+    note.source = source;
+    note.message = note_msg;
+    add(std::move(note));
   }
 
   /// adds the warning message with the given Source to the end of this list.
+  /// @param system the system raising the warning message
   /// @param warning_msg the warning message
   /// @param source the source of the warning diagnostic
-  void add_warning(const std::string& warning_msg, const Source& source) {
-    diag::Diagnostic error{};
-    error.severity = diag::Severity::Warning;
-    error.source = source;
-    error.message = warning_msg;
-    add(std::move(error));
+  void add_warning(System system,
+                   const std::string& warning_msg,
+                   const Source& source) {
+    diag::Diagnostic warning{};
+    warning.severity = diag::Severity::Warning;
+    warning.system = system;
+    warning.source = source;
+    warning.message = warning_msg;
+    add(std::move(warning));
   }
 
   /// adds the error message without a source to the end of this list.
+  /// @param system the system raising the error message
   /// @param err_msg the error message
-  void add_error(std::string err_msg) {
+  void add_error(System system, std::string err_msg) {
     diag::Diagnostic error{};
     error.severity = diag::Severity::Error;
+    error.system = system;
     error.message = std::move(err_msg);
     add(std::move(error));
   }
 
   /// adds the error message with the given Source to the end of this list.
+  /// @param system the system raising the error message
   /// @param err_msg the error message
   /// @param source the source of the error diagnostic
-  void add_error(std::string err_msg, const Source& source) {
+  void add_error(System system, std::string err_msg, const Source& source) {
     diag::Diagnostic error{};
     error.severity = diag::Severity::Error;
+    error.system = system;
     error.source = source;
     error.message = std::move(err_msg);
     add(std::move(error));
@@ -141,24 +173,33 @@
 
   /// adds the error message with the given code and Source to the end of this
   /// list.
+  /// @param system the system raising the error message
   /// @param code the error code
   /// @param err_msg the error message
   /// @param source the source of the error diagnostic
-  void add_error(const char* code, std::string err_msg, const Source& source) {
+  void add_error(System system,
+                 const char* code,
+                 std::string err_msg,
+                 const Source& source) {
     diag::Diagnostic error{};
     error.code = code;
     error.severity = diag::Severity::Error;
+    error.system = system;
     error.source = source;
     error.message = std::move(err_msg);
     add(std::move(error));
   }
 
   /// adds an internal compiler error message to the end of this list.
+  /// @param system the system raising the error message
   /// @param err_msg the error message
   /// @param source the source of the internal compiler error
-  void add_ice(const std::string& err_msg, const Source& source) {
+  void add_ice(System system,
+               const std::string& err_msg,
+               const Source& source) {
     diag::Diagnostic ice{};
     ice.severity = diag::Severity::InternalCompilerError;
+    ice.system = system;
     ice.source = source;
     ice.message = err_msg;
     add(std::move(ice));
diff --git a/src/diagnostic/formatter_test.cc b/src/diagnostic/formatter_test.cc
index 6f62243..08565ab 100644
--- a/src/diagnostic/formatter_test.cc
+++ b/src/diagnostic/formatter_test.cc
@@ -33,18 +33,19 @@
   Source::File file{"file.name", content};
   Diagnostic diag_note{Severity::Note,
                        Source{Source::Range{Source::Location{1, 14}}, &file},
-                       "purr"};
+                       "purr", System::Test};
   Diagnostic diag_warn{Severity::Warning,
-                       Source{Source::Range{{2, 14}, {2, 18}}, &file}, "grrr"};
+                       Source{Source::Range{{2, 14}, {2, 18}}, &file}, "grrr",
+                       System::Test};
   Diagnostic diag_err{Severity::Error,
                       Source{Source::Range{{3, 16}, {3, 21}}, &file}, "hiss",
-                      "abc123"};
+                      System::Test, "abc123"};
   Diagnostic diag_ice{Severity::InternalCompilerError,
                       Source{Source::Range{{4, 16}, {4, 19}}, &file},
-                      "unreachable"};
+                      "unreachable", System::Test};
   Diagnostic diag_fatal{Severity::Fatal,
                         Source{Source::Range{{4, 16}, {4, 19}}, &file},
-                        "nothing"};
+                        "nothing", System::Test};
 };
 
 TEST_F(DiagFormatterTest, Simple) {
@@ -68,7 +69,7 @@
 
 TEST_F(DiagFormatterTest, SimpleNoSource) {
   Formatter fmt{{false, false, false, false}};
-  Diagnostic diag{Severity::Note, Source{}, "no source!"};
+  Diagnostic diag{Severity::Note, Source{}, "no source!", System::Test};
   auto got = fmt.format(List{diag});
   auto* expect = "no source!";
   ASSERT_EQ(expect, got);
@@ -131,7 +132,7 @@
 TEST_F(DiagFormatterTest, BasicWithMultiLine) {
   Diagnostic multiline{Severity::Warning,
                        Source{Source::Range{{2, 9}, {4, 15}}, &file},
-                       "multiline"};
+                       "multiline", System::Test};
   Formatter fmt{{false, false, true, false}};
   auto got = fmt.format(List{multiline});
   auto* expect = R"(2:9: multiline
@@ -166,7 +167,7 @@
 TEST_F(DiagFormatterTest, BasicWithMultiLineTab4) {
   Diagnostic multiline{Severity::Warning,
                        Source{Source::Range{{2, 9}, {4, 15}}, &file},
-                       "multiline"};
+                       "multiline", System::Test};
   Formatter fmt{{false, false, true, false, 4u}};
   auto got = fmt.format(List{multiline});
   auto* expect = R"(2:9: multiline
@@ -219,7 +220,7 @@
 TEST_F(DiagFormatterTest, RangeOOB) {
   Formatter fmt{{true, true, true, true}};
   diag::List list;
-  list.add_error("oob", Source{{{10, 20}, {30, 20}}, &file});
+  list.add_error(System::Test, "oob", Source{{{10, 20}, {30, 20}}, &file});
   auto got = fmt.format(list);
   auto* expect = R"(file.name:10:20 error: oob
 
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index a757d55..c99e6b8 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -47,7 +47,7 @@
 
 void AppendResourceBindings(std::vector<ResourceBinding>* dest,
                             const std::vector<ResourceBinding>& orig) {
-  TINT_ASSERT(dest);
+  TINT_ASSERT(Inspector, dest);
   if (!dest) {
     return;
   }
@@ -84,7 +84,7 @@
     if (wgsize[0].overridable_const || wgsize[1].overridable_const ||
         wgsize[2].overridable_const) {
       // TODO(crbug.com/tint/713): Handle overridable constants.
-      TINT_ASSERT(false);
+      TINT_ASSERT(Inspector, false);
     }
 
     for (auto* param : sem->Parameters()) {
@@ -528,7 +528,7 @@
   }
 
   auto* location = ast::GetDecoration<ast::LocationDecoration>(decorations);
-  TINT_ASSERT(location != nullptr);
+  TINT_ASSERT(Inspector, location != nullptr);
   stage_variable.has_location_decoration = true;
   stage_variable.location_decoration = location->value();
 
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc
index 4b73992..0997eae 100644
--- a/src/intrinsic_table.cc
+++ b/src/intrinsic_table.cc
@@ -120,11 +120,11 @@
   const sem::Type* Type(uint32_t idx) const {
     auto it = types_.find(idx);
     if (it == types_.end()) {
-      TINT_ICE(builder.Diagnostics())
+      TINT_ICE(Resolver, builder.Diagnostics())
           << "type with index " << idx << " is not closed";
       return nullptr;
     }
-    TINT_ASSERT(it != types_.end());
+    TINT_ASSERT(Resolver, it != types_.end());
     return it->second;
   }
 
@@ -133,7 +133,7 @@
   Number Num(uint32_t idx) const {
     auto it = numbers_.find(idx);
     if (it == numbers_.end()) {
-      TINT_ICE(builder.Diagnostics())
+      TINT_ICE(Resolver, builder.Diagnostics())
           << "number with index " << idx << " is not closed";
       return Number::invalid;
     }
@@ -802,7 +802,7 @@
       ss << std::endl;
     }
   }
-  builder.Diagnostics().add_error(ss.str(), source);
+  builder.Diagnostics().add_error(diag::System::Resolver, ss.str(), source);
   return nullptr;
 }
 
@@ -888,7 +888,7 @@
     if (!return_type) {
       std::stringstream ss;
       PrintOverload(ss, overload, intrinsic_type);
-      TINT_ICE(builder.Diagnostics())
+      TINT_ICE(Resolver, builder.Diagnostics())
           << "MatchState.Match() returned null for " << ss.str();
       return nullptr;
     }
diff --git a/src/program.cc b/src/program.cc
index 7a0135b..81c24f0 100644
--- a/src/program.cc
+++ b/src/program.cc
@@ -63,7 +63,7 @@
     // If the builder claims to be invalid, then we really should have an error
     // message generated. If we find a situation where the program is not valid
     // and there are no errors reported, add one here.
-    diagnostics_.add_error("invalid program generated");
+    diagnostics_.add_error(diag::System::Program, "invalid program generated");
   }
 }
 
@@ -128,7 +128,7 @@
 }
 
 void Program::AssertNotMoved() const {
-  TINT_ASSERT(!moved_);
+  TINT_ASSERT(Program, !moved_);
 }
 
 }  // namespace tint
diff --git a/src/program_builder.cc b/src/program_builder.cc
index 632d480..effe513 100644
--- a/src/program_builder.cc
+++ b/src/program_builder.cc
@@ -83,7 +83,7 @@
 
 void ProgramBuilder::AssertNotMoved() const {
   if (moved_) {
-    TINT_ICE(const_cast<ProgramBuilder*>(this)->diagnostics_)
+    TINT_ICE(ProgramBuilder, const_cast<ProgramBuilder*>(this)->diagnostics_)
         << "Attempting to use ProgramBuilder after it has been moved";
   }
 }
diff --git a/src/program_id.h b/src/program_id.h
index f6e0f36..bd98220 100644
--- a/src/program_id.h
+++ b/src/program_id.h
@@ -89,6 +89,7 @@
 void AssertProgramIDsEqual(A&& a,
                            B&& b,
                            bool if_valid,
+                           diag::System system,
                            const char* msg,
                            const char* file,
                            size_t line) {
@@ -101,25 +102,27 @@
     return;  //  a or b were not valid
   }
   diag::List diagnostics;
-  tint::InternalCompilerError(file, line, diagnostics) << msg;
+  tint::InternalCompilerError(file, line, system, diagnostics) << msg;
 }
 
 }  // namespace detail
 
-/// TINT_ASSERT_PROGRAM_IDS_EQUAL(A, B) is a macro that asserts that the program
-/// identifiers for A and B are equal.
+/// TINT_ASSERT_PROGRAM_IDS_EQUAL(SYSTEM, A, B) is a macro that asserts that the
+/// program identifiers for A and B are equal.
 ///
-/// TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(A, B) is a macro that asserts that
-/// the program identifiers for A and B are equal, if both A and B have valid
-/// program identifiers.
+/// TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(SYSTEM, A, B) is a macro that asserts
+/// that the program identifiers for A and B are equal, if both A and B have
+/// valid program identifiers.
 #if TINT_CHECK_FOR_CROSS_PROGRAM_LEAKS
-#define TINT_ASSERT_PROGRAM_IDS_EQUAL(a, b)                                   \
-  detail::AssertProgramIDsEqual(                                              \
-      a, b, false, "TINT_ASSERT_PROGRAM_IDS_EQUAL(" #a ", " #b ")", __FILE__, \
-      __LINE__)
-#define TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(a, b)                        \
-  detail::AssertProgramIDsEqual(                                            \
-      a, b, true, "TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(" #a ", " #b ")", \
+#define TINT_ASSERT_PROGRAM_IDS_EQUAL(system, a, b)                      \
+  detail::AssertProgramIDsEqual(a, b, false, tint::diag::System::system, \
+                                "TINT_ASSERT_PROGRAM_IDS_EQUAL(" #system \
+                                "," #a ", " #b ")",                      \
+                                __FILE__, __LINE__)
+#define TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(system, a, b)                 \
+  detail::AssertProgramIDsEqual(                                             \
+      a, b, true, tint::diag::System::system,                                \
+      "TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(" #system ", " #a ", " #b ")", \
       __FILE__, __LINE__)
 #else
 #define TINT_ASSERT_PROGRAM_IDS_EQUAL(a, b) \
diff --git a/src/program_test.cc b/src/program_test.cc
index acd6355..1f60df8 100644
--- a/src/program_test.cc
+++ b/src/program_test.cc
@@ -87,7 +87,7 @@
 }
 
 TEST_F(ProgramTest, DiagnosticsMove) {
-  Diagnostics().add_error("an error message");
+  Diagnostics().add_error(diag::System::Program, "an error message");
 
   Program program_a(std::move(*this));
   EXPECT_FALSE(program_a.IsValid());
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index 84706d6..c8b4df0 100644
--- a/src/reader/spirv/function.cc
+++ b/src/reader/spirv/function.cc
@@ -779,7 +779,7 @@
 FunctionEmitter::StatementBlock::~StatementBlock() = default;
 
 void FunctionEmitter::StatementBlock::Finalize(ProgramBuilder* pb) {
-  TINT_ASSERT(!finalized_ /* Finalize() must only be called once */);
+  TINT_ASSERT(Reader, !finalized_ /* Finalize() must only be called once */);
 
   for (size_t i = 0; i < statements_.size(); i++) {
     if (auto* sb = statements_[i]->As<StatementBuilder>()) {
@@ -795,7 +795,8 @@
 }
 
 void FunctionEmitter::StatementBlock::Add(ast::Statement* statement) {
-  TINT_ASSERT(!finalized_ /* Add() must not be called after Finalize() */);
+  TINT_ASSERT(Reader,
+              !finalized_ /* Add() must not be called after Finalize() */);
   statements_.emplace_back(statement);
 }
 
@@ -807,8 +808,8 @@
 
 void FunctionEmitter::PushGuard(const std::string& guard_name,
                                 uint32_t end_id) {
-  TINT_ASSERT(!statements_stack_.empty());
-  TINT_ASSERT(!guard_name.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
+  TINT_ASSERT(Reader, !guard_name.empty());
   // Guard control flow by the guard variable.  Introduce a new
   // if-selection with a then-clause ending at the same block
   // as the statement block at the top of the stack.
@@ -825,7 +826,7 @@
 }
 
 void FunctionEmitter::PushTrueGuard(uint32_t end_id) {
-  TINT_ASSERT(!statements_stack_.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
   const auto& top = statements_stack_.back();
 
   auto* cond = MakeTrue(Source{});
@@ -838,14 +839,14 @@
 }
 
 const ast::StatementList FunctionEmitter::ast_body() {
-  TINT_ASSERT(!statements_stack_.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
   auto& entry = statements_stack_[0];
   entry.Finalize(&builder_);
   return entry.GetStatements();
 }
 
 ast::Statement* FunctionEmitter::AddStatement(ast::Statement* statement) {
-  TINT_ASSERT(!statements_stack_.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
   if (statement != nullptr) {
     statements_stack_.back().Add(statement);
   }
@@ -853,9 +854,9 @@
 }
 
 ast::Statement* FunctionEmitter::LastStatement() {
-  TINT_ASSERT(!statements_stack_.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
   auto& statement_list = statements_stack_.back().GetStatements();
-  TINT_ASSERT(!statement_list.empty());
+  TINT_ASSERT(Reader, !statement_list.empty());
   return statement_list.back();
 }
 
@@ -877,7 +878,7 @@
 
   bool make_body_function = true;
   if (ep_info_) {
-    TINT_ASSERT(!ep_info_->inner_name.empty());
+    TINT_ASSERT(Reader, !ep_info_->inner_name.empty());
     if (ep_info_->owns_inner_implementation) {
       // This is an entry point, and we want to emit it as a wrapper around
       // an implementation function.
@@ -909,7 +910,7 @@
 }
 
 ast::BlockStatement* FunctionEmitter::MakeFunctionBody() {
-  TINT_ASSERT(statements_stack_.size() == 1);
+  TINT_ASSERT(Reader, statements_stack_.size() == 1);
 
   if (!EmitBody()) {
     return nullptr;
@@ -950,8 +951,8 @@
   // have already been created.
   for (uint32_t var_id : ep_info_->inputs) {
     const auto* var = def_use_mgr_->GetDef(var_id);
-    TINT_ASSERT(var != nullptr);
-    TINT_ASSERT(var->opcode() == SpvOpVariable);
+    TINT_ASSERT(Reader, var != nullptr);
+    TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
     auto* store_type = GetVariableStoreType(*var);
     auto* forced_store_type = store_type;
     ast::DecorationList param_decos;
@@ -1050,8 +1051,8 @@
             create<ast::BuiltinDecoration>(source, ast::Builtin::kPosition));
       } else {
         const auto* var = def_use_mgr_->GetDef(var_id);
-        TINT_ASSERT(var != nullptr);
-        TINT_ASSERT(var->opcode() == SpvOpVariable);
+        TINT_ASSERT(Reader, var != nullptr);
+        TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
         store_type = GetVariableStoreType(*var);
         param_type = store_type;
         if (!parser_impl_.ConvertDecorationsForVariable(var_id, &param_type,
@@ -1502,7 +1503,7 @@
   //      block. Also mark the the most recent continue target for which we
   //      haven't reached the backedge block.
 
-  TINT_ASSERT(block_order_.size() > 0);
+  TINT_ASSERT(Reader, block_order_.size() > 0);
   constructs_.clear();
   const auto entry_id = block_order_[0];
 
@@ -1523,8 +1524,8 @@
     // A loop construct is added right after its associated continue construct.
     // In that case, adjust the parent up.
     if (k == Construct::kLoop) {
-      TINT_ASSERT(parent);
-      TINT_ASSERT(parent->kind == Construct::kContinue);
+      TINT_ASSERT(Reader, parent);
+      TINT_ASSERT(Reader, parent->kind == Construct::kContinue);
       scope_end_pos = parent->end_pos;
       parent = parent->parent;
     }
@@ -1543,9 +1544,9 @@
 
   for (uint32_t i = 0; i < block_order_.size(); ++i) {
     const auto block_id = block_order_[i];
-    TINT_ASSERT(block_id > 0);
+    TINT_ASSERT(Reader, block_id > 0);
     auto* block_info = GetBlockInfo(block_id);
-    TINT_ASSERT(block_info);
+    TINT_ASSERT(Reader, block_info);
 
     if (enclosing.empty()) {
       return Fail() << "internal error: too many merge blocks before block "
@@ -1619,7 +1620,7 @@
       }
     }
 
-    TINT_ASSERT(top);
+    TINT_ASSERT(Reader, top);
     block_info->construct = top;
   }
 
@@ -1828,9 +1829,9 @@
   //    NEC(S) is the parent of NEC(T).
 
   for (const auto src : block_order_) {
-    TINT_ASSERT(src > 0);
+    TINT_ASSERT(Reader, src > 0);
     auto* src_info = GetBlockInfo(src);
-    TINT_ASSERT(src_info);
+    TINT_ASSERT(Reader, src_info);
     const auto src_pos = src_info->pos;
     const auto& src_construct = *(src_info->construct);
 
@@ -1868,7 +1869,7 @@
     for (const auto dest : successors) {
       const auto* dest_info = GetBlockInfo(dest);
       // We've already checked terminators are valid.
-      TINT_ASSERT(dest_info);
+      TINT_ASSERT(Reader, dest_info);
       const auto dest_pos = dest_info->pos;
 
       // Insert the edge kind entry and keep a handle to update
@@ -1893,7 +1894,7 @@
                         << " (violates post-dominance rule)";
         }
         const auto* ct_info = GetBlockInfo(continue_construct->begin_id);
-        TINT_ASSERT(ct_info);
+        TINT_ASSERT(Reader, ct_info);
         if (ct_info->header_for_continue != dest) {
           return Fail()
                  << "Invalid backedge (" << src << "->" << dest
@@ -2185,7 +2186,7 @@
       // The first clause might be a then-clause or an else-clause.
       const auto second_head = std::max(true_head_pos, false_head_pos);
       const auto end_first_clause_pos = second_head - 1;
-      TINT_ASSERT(end_first_clause_pos < block_order_.size());
+      TINT_ASSERT(Reader, end_first_clause_pos < block_order_.size());
       const auto end_first_clause = block_order_[end_first_clause_pos];
       uint32_t premerge_id = 0;
       uint32_t if_break_id = 0;
@@ -2405,15 +2406,15 @@
 
   // Upon entry, the statement stack has one entry representing the whole
   // function.
-  TINT_ASSERT(!constructs_.empty());
+  TINT_ASSERT(Reader, !constructs_.empty());
   Construct* function_construct = constructs_[0].get();
-  TINT_ASSERT(function_construct != nullptr);
-  TINT_ASSERT(function_construct->kind == Construct::kFunction);
+  TINT_ASSERT(Reader, function_construct != nullptr);
+  TINT_ASSERT(Reader, function_construct->kind == Construct::kFunction);
   // Make the first entry valid by filling in the construct field, which
   // had not been computed at the time the entry was first created.
   // TODO(dneto): refactor how the first construct is created vs.
   // this statements stack entry is populated.
-  TINT_ASSERT(statements_stack_.size() == 1);
+  TINT_ASSERT(Reader, statements_stack_.size() == 1);
   statements_stack_[0].SetConstruct(function_construct);
 
   for (auto block_id : block_order()) {
@@ -2612,8 +2613,8 @@
 bool FunctionEmitter::EmitIfStart(const BlockInfo& block_info) {
   // The block is the if-header block.  So its construct is the if construct.
   auto* construct = block_info.construct;
-  TINT_ASSERT(construct->kind == Construct::kIfSelection);
-  TINT_ASSERT(construct->begin_id == block_info.id);
+  TINT_ASSERT(Reader, construct->kind == Construct::kIfSelection);
+  TINT_ASSERT(Reader, construct->begin_id == block_info.id);
 
   const uint32_t true_head = block_info.true_head;
   const uint32_t false_head = block_info.false_head;
@@ -2757,8 +2758,8 @@
 bool FunctionEmitter::EmitSwitchStart(const BlockInfo& block_info) {
   // The block is the if-header block.  So its construct is the if construct.
   auto* construct = block_info.construct;
-  TINT_ASSERT(construct->kind == Construct::kSwitchSelection);
-  TINT_ASSERT(construct->begin_id == block_info.id);
+  TINT_ASSERT(Reader, construct->kind == Construct::kSwitchSelection);
+  TINT_ASSERT(Reader, construct->begin_id == block_info.id);
   const auto* branch = block_info.basic_block->terminator();
 
   const auto selector_id = branch->GetSingleWordInOperand(0);
@@ -2806,7 +2807,7 @@
       clause_heads[w] = clause_heads[r];
     }
     // We know it's not empty because it always has at least a default clause.
-    TINT_ASSERT(!clause_heads.empty());
+    TINT_ASSERT(Reader, !clause_heads.empty());
     clause_heads.resize(w + 1);
   }
 
@@ -3019,9 +3020,10 @@
       // Unless forced, don't bother with a break at the end of a case/default
       // clause.
       const auto header = dest_info.header_for_merge;
-      TINT_ASSERT(header != 0);
+      TINT_ASSERT(Reader, header != 0);
       const auto* exiting_construct = GetBlockInfo(header)->construct;
-      TINT_ASSERT(exiting_construct->kind == Construct::kSwitchSelection);
+      TINT_ASSERT(Reader,
+                  exiting_construct->kind == Construct::kSwitchSelection);
       const auto candidate_next_case_pos = src_info.pos + 1;
       // Leaving the last block from the last case?
       if (candidate_next_case_pos == dest_info.pos) {
@@ -3165,7 +3167,7 @@
   // Emit declarations of hoisted variables, in index order.
   for (auto id : sorted_by_index(block_info.hoisted_ids)) {
     const auto* def_inst = def_use_mgr_->GetDef(id);
-    TINT_ASSERT(def_inst);
+    TINT_ASSERT(Reader, def_inst);
     auto* storage_type =
         RemapStorageClass(parser_impl_.ConvertType(def_inst->type_id()), id);
     AddStatement(create<ast::VariableDeclStatement>(
@@ -3178,9 +3180,9 @@
   // Emit declarations of phi state variables, in index order.
   for (auto id : sorted_by_index(block_info.phis_needing_state_vars)) {
     const auto* def_inst = def_use_mgr_->GetDef(id);
-    TINT_ASSERT(def_inst);
+    TINT_ASSERT(Reader, def_inst);
     const auto phi_var_name = GetDefInfo(id)->phi_var;
-    TINT_ASSERT(!phi_var_name.empty());
+    TINT_ASSERT(Reader, !phi_var_name.empty());
     auto* var = builder_.Var(
         phi_var_name,
         parser_impl_.ConvertType(def_inst->type_id())->Build(builder_));
@@ -3382,7 +3384,7 @@
           if (auto* arr = lhs.type->As<Array>()) {
             lhs.type = arr->type;
           }
-          TINT_ASSERT(lhs.type);
+          TINT_ASSERT(Reader, lhs.type);
           break;
         default:
           break;
@@ -4008,7 +4010,7 @@
     const auto pointer_type_id =
         type_mgr_->FindPointerToType(pointee_type_id, storage_class);
     auto* type = parser_impl_.ConvertType(pointer_type_id, PtrAs::Ref);
-    TINT_ASSERT(type && type->Is<Reference>());
+    TINT_ASSERT(Reader, type && type->Is<Reference>());
     current_expr = TypedExpression{type, next_expr};
   }
   return current_expr;
@@ -4206,7 +4208,7 @@
           source, expr.expr, Swizzle(index)));
     } else if (index < vec0_len + vec1_len) {
       const auto sub_index = index - vec0_len;
-      TINT_ASSERT(sub_index < kMaxVectorLen);
+      TINT_ASSERT(Reader, sub_index < kMaxVectorLen);
       auto expr = MakeExpression(vec1_id);
       if (!expr) {
         return {};
@@ -4511,7 +4513,7 @@
                                                     uint32_t last_pos) const {
   const auto* enclosing_construct =
       GetBlockInfo(block_order_[first_pos])->construct;
-  TINT_ASSERT(enclosing_construct != nullptr);
+  TINT_ASSERT(Reader, enclosing_construct != nullptr);
   // Constructs are strictly nesting, so follow parent pointers
   while (enclosing_construct &&
          !enclosing_construct->ScopeContainsPos(last_pos)) {
@@ -4523,7 +4525,7 @@
         sibling_loop ? sibling_loop : enclosing_construct->parent;
   }
   // At worst, we go all the way out to the function construct.
-  TINT_ASSERT(enclosing_construct != nullptr);
+  TINT_ASSERT(Reader, enclosing_construct != nullptr);
   return enclosing_construct;
 }
 
@@ -5270,11 +5272,11 @@
 
   // The texel type is always a 4-element vector.
   const uint32_t dest_count = 4u;
-  TINT_ASSERT(dest_type->Is<Vector>() &&
-              dest_type->As<Vector>()->size == dest_count);
-  TINT_ASSERT(dest_type->IsFloatVector() ||
-              dest_type->IsUnsignedIntegerVector() ||
-              dest_type->IsSignedIntegerVector());
+  TINT_ASSERT(Reader, dest_type->Is<Vector>() &&
+                          dest_type->As<Vector>()->size == dest_count);
+  TINT_ASSERT(Reader, dest_type->IsFloatVector() ||
+                          dest_type->IsUnsignedIntegerVector() ||
+                          dest_type->IsSignedIntegerVector());
 
   if (src_type == dest_type) {
     return texel.expr;
@@ -5294,7 +5296,7 @@
   }
 
   const auto required_count = parser_impl_.GetChannelCountForFormat(format);
-  TINT_ASSERT(0 < required_count && required_count <= 4);
+  TINT_ASSERT(Reader, 0 < required_count && required_count <= 4);
 
   const uint32_t src_count =
       src_type->IsScalar() ? 1 : src_type->As<Vector>()->size;
diff --git a/src/reader/spirv/function.h b/src/reader/spirv/function.h
index ef920ba..bdbe5e2 100644
--- a/src/reader/spirv/function.h
+++ b/src/reader/spirv/function.h
@@ -1017,7 +1017,7 @@
   /// @return the built StatementBuilder
   template <typename T, typename... ARGS>
   T* AddStatementBuilder(ARGS&&... args) {
-    TINT_ASSERT(!statements_stack_.empty());
+    TINT_ASSERT(Reader, !statements_stack_.empty());
     return statements_stack_.back().AddStatementBuilder<T>(
         std::forward<ARGS>(args)...);
   }
diff --git a/src/reader/spirv/namer.cc b/src/reader/spirv/namer.cc
index b459712..4d819b4 100644
--- a/src/reader/spirv/namer.cc
+++ b/src/reader/spirv/namer.cc
@@ -123,7 +123,7 @@
 std::string Namer::MakeDerivedName(const std::string& base_name) {
   auto result = FindUnusedDerivedName(base_name);
   const bool registered = RegisterWithoutId(result);
-  TINT_ASSERT(registered);
+  TINT_ASSERT(Reader, registered);
   return result;
 }
 
diff --git a/src/reader/spirv/parser.cc b/src/reader/spirv/parser.cc
index 24000bf..c382d90 100644
--- a/src/reader/spirv/parser.cc
+++ b/src/reader/spirv/parser.cc
@@ -29,7 +29,7 @@
   ProgramBuilder& builder = parser.builder();
   if (!parsed) {
     // TODO(bclayton): Migrate spirv::ParserImpl to using diagnostics.
-    builder.Diagnostics().add_error(parser.error());
+    builder.Diagnostics().add_error(diag::System::Reader, parser.error());
     return Program(std::move(builder));
   }
 
diff --git a/src/reader/spirv/parser_impl.cc b/src/reader/spirv/parser_impl.cc
index 94fe563..5272aa1 100644
--- a/src/reader/spirv/parser_impl.cc
+++ b/src/reader/spirv/parser_impl.cc
@@ -833,8 +833,8 @@
       // Reuse the inner implementation owned by the first entry point.
       inner_implementation_name = where->second[0].inner_name;
     }
-    TINT_ASSERT(!inner_implementation_name.empty());
-    TINT_ASSERT(ep_name != inner_implementation_name);
+    TINT_ASSERT(Reader, !inner_implementation_name.empty());
+    TINT_ASSERT(Reader, ep_name != inner_implementation_name);
 
     tint::UniqueVector<uint32_t> inputs;
     tint::UniqueVector<uint32_t> outputs;
diff --git a/src/reader/spirv/parser_impl_barrier_test.cc b/src/reader/spirv/parser_impl_barrier_test.cc
index 0f82e00..303f5fe 100644
--- a/src/reader/spirv/parser_impl_barrier_test.cc
+++ b/src/reader/spirv/parser_impl_barrier_test.cc
@@ -40,7 +40,7 @@
   auto p = std::make_unique<ParserImpl>(test::Assemble(preamble + spirv));
   if (!p->BuildAndParseInternalModule()) {
     ProgramBuilder builder;
-    builder.Diagnostics().add_error(p->error());
+    builder.Diagnostics().add_error(diag::System::Reader, p->error());
     return Program(std::move(builder));
   }
   return p->program();
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index 71ba2d0..3930299 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -269,14 +269,15 @@
 ParserImpl::Failure::Errored ParserImpl::add_error(const Source& source,
                                                    const std::string& err) {
   if (silence_errors_ == 0) {
-    builder_.Diagnostics().add_error(err, source);
+    builder_.Diagnostics().add_error(diag::System::Reader, err, source);
   }
   return Failure::kErrored;
 }
 
 void ParserImpl::deprecated(const Source& source, const std::string& msg) {
   builder_.Diagnostics().add_warning(
-      "use of deprecated language feature: " + msg, source);
+      diag::System::Reader, "use of deprecated language feature: " + msg,
+      source);
 }
 
 Token ParserImpl::next() {
@@ -3298,7 +3299,7 @@
   --parse_depth_;
 
   if (sync_tokens_.back() != tok) {
-    TINT_ICE(builder_.Diagnostics()) << "sync_tokens is out of sync";
+    TINT_ICE(Reader, builder_.Diagnostics()) << "sync_tokens is out of sync";
   }
   sync_tokens_.pop_back();
 
diff --git a/src/reader/wgsl/parser_impl.h b/src/reader/wgsl/parser_impl.h
index 55cd25a..19b29b8 100644
--- a/src/reader/wgsl/parser_impl.h
+++ b/src/reader/wgsl/parser_impl.h
@@ -112,7 +112,7 @@
     /// return type will always be a pointer to a non-pointer type. #errored
     /// must be false to call.
     inline typename detail::OperatorArrow<T>::type operator->() {
-      TINT_ASSERT(!errored);
+      TINT_ASSERT(Reader, !errored);
       return detail::OperatorArrow<T>::ptr(value);
     }
 
@@ -183,7 +183,7 @@
     /// return type will always be a pointer to a non-pointer type. #errored
     /// must be false to call.
     inline typename detail::OperatorArrow<T>::type operator->() {
-      TINT_ASSERT(!errored);
+      TINT_ASSERT(Reader, !errored);
       return detail::OperatorArrow<T>::ptr(value);
     }
 
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 8761eaa..9889a08 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -164,7 +164,8 @@
   bool result = ResolveInternal();
 
   if (!result && !diagnostics_.contains_errors()) {
-    TINT_ICE(diagnostics_) << "resolving failed, but no error was raised";
+    TINT_ICE(Resolver, diagnostics_)
+        << "resolving failed, but no error was raised";
     return false;
   }
 
@@ -263,7 +264,7 @@
         return false;
       }
     } else {
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Resolver, diagnostics_)
           << "unhandled global declaration: " << decl->TypeInfo().name;
       return false;
     }
@@ -280,11 +281,12 @@
 
   for (auto* node : builder_->ASTNodes().Objects()) {
     if (marked_.count(node) == 0) {
-      TINT_ICE(diagnostics_) << "AST node '" << node->TypeInfo().name
-                             << "' was not reached by the resolver\n"
-                             << "At: " << node->source() << "\n"
-                             << "Content: " << builder_->str(node) << "\n"
-                             << "Pointer: " << node;
+      TINT_ICE(Resolver, diagnostics_)
+          << "AST node '" << node->TypeInfo().name
+          << "' was not reached by the resolver\n"
+          << "At: " << node->source() << "\n"
+          << "Content: " << builder_->str(node) << "\n"
+          << "Pointer: " << node;
       result = false;
     }
   }
@@ -392,14 +394,14 @@
     if (auto* t = ty->As<ast::TypeName>()) {
       auto it = named_type_info_.find(t->name());
       if (it == named_type_info_.end()) {
-        diagnostics_.add_error(
+        AddError(
             "unknown type '" + builder_->Symbols().NameFor(t->name()) + "'",
             t->source());
         return nullptr;
       }
       return it->second.sem;
     }
-    TINT_UNREACHABLE(diagnostics_)
+    TINT_UNREACHABLE(Resolver, diagnostics_)
         << "Unhandled ast::Type: " << ty->TypeInfo().name;
     return nullptr;
   }();
@@ -414,8 +416,7 @@
   // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types
   // T must be either u32 or i32.
   if (!s->Type()->IsAnyOf<sem::U32, sem::I32>()) {
-    diagnostics_.add_error("atomic only supports i32 or u32 types",
-                           a->type()->source());
+    AddError("atomic only supports i32 or u32 types", a->type()->source());
     return false;
   }
   return true;
@@ -435,7 +436,7 @@
         // TODO(crbug.com/tint/901): Validate that the access mode is
         // read_write.
         auto* member = structure->Members()[sm.index];
-        diagnostics_.add_error(
+        AddError(
             "atomic types can only be used in storage classes workgroup or "
             "storage, but was used by storage class " +
                 std::string(ast::str(usage)),
@@ -452,13 +453,11 @@
 bool Resolver::ValidateStorageTexture(const ast::StorageTexture* t) {
   switch (t->access()) {
     case ast::Access::kUndefined:
-      diagnostics_.add_error("storage textures must have access control",
-                             t->source());
+      AddError("storage textures must have access control", t->source());
       return false;
     case ast::Access::kReadWrite:
-      diagnostics_.add_error(
-          "storage textures only support read-only and write-only access",
-          t->source());
+      AddError("storage textures only support read-only and write-only access",
+               t->source());
       return false;
 
     case ast::Access::kRead:
@@ -467,13 +466,13 @@
   }
 
   if (!IsValidStorageTextureDimension(t->dim())) {
-    diagnostics_.add_error(
-        "cube dimensions for storage textures are not supported", t->source());
+    AddError("cube dimensions for storage textures are not supported",
+             t->source());
     return false;
   }
 
   if (!IsValidStorageTextureImageFormat(t->image_format())) {
-    diagnostics_.add_error(
+    AddError(
         "image format must be one of the texel formats specified for storage "
         "textues in https://gpuweb.github.io/gpuweb/wgsl/#texel-formats",
         t->source());
@@ -485,9 +484,9 @@
 Resolver::VariableInfo* Resolver::Variable(ast::Variable* var,
                                            VariableKind kind) {
   if (variable_to_info_.count(var)) {
-    TINT_ICE(diagnostics_) << "Variable "
-                           << builder_->Symbols().NameFor(var->symbol())
-                           << " already resolved";
+    TINT_ICE(Resolver, diagnostics_)
+        << "Variable " << builder_->Symbols().NameFor(var->symbol())
+        << " already resolved";
     return nullptr;
   }
 
@@ -523,8 +522,7 @@
     // If the variable has no declared type, infer it from the RHS
     if (!storage_type) {
       if (!var->is_const() && kind == VariableKind::kGlobal) {
-        diagnostics_.add_error("global var declaration must specify a type",
-                               var->source());
+        AddError("global var declaration must specify a type", var->source());
         return nullptr;
       }
 
@@ -533,13 +531,12 @@
     }
   } else if (var->is_const() && kind != VariableKind::kParameter &&
              !ast::HasDecoration<ast::OverrideDecoration>(var->decorations())) {
-    diagnostics_.add_error("let declarations must have initializers",
-                           var->source());
+    AddError("let declarations must have initializers", var->source());
     return nullptr;
   }
 
   if (!storage_type) {
-    TINT_ICE(diagnostics_)
+    TINT_ICE(Resolver, diagnostics_)
         << "failed to determine storage type for variable '" +
                builder_->Symbols().NameFor(var->symbol()) + "'\n"
         << "Source: " << var->source();
@@ -608,10 +605,9 @@
   // Value type has to match storage type
   if (storage_type != value_type) {
     std::string decl = var->is_const() ? "let" : "var";
-    diagnostics_.add_error("cannot initialize " + decl + " of type '" +
-                               type_name + "' with value of type '" +
-                               rhs_type_name + "'",
-                           var->source());
+    AddError("cannot initialize " + decl + " of type '" + type_name +
+                 "' with value of type '" + rhs_type_name + "'",
+             var->source());
     return false;
   }
   return true;
@@ -619,10 +615,10 @@
 
 bool Resolver::GlobalVariable(ast::Variable* var) {
   if (variable_stack_.has(var->symbol())) {
-    diagnostics_.add_error("v-0011",
-                           "redeclared global identifier '" +
-                               builder_->Symbols().NameFor(var->symbol()) + "'",
-                           var->source());
+    AddError("v-0011",
+             "redeclared global identifier '" +
+                 builder_->Symbols().NameFor(var->symbol()) + "'",
+             var->source());
     return false;
   }
 
@@ -633,14 +629,13 @@
   variable_stack_.set_global(var->symbol(), info);
 
   if (!var->is_const() && info->storage_class == ast::StorageClass::kNone) {
-    diagnostics_.add_error(
-        "v-0022", "global variables must have a storage class", var->source());
+    AddError("v-0022", "global variables must have a storage class",
+             var->source());
     return false;
   }
   if (var->is_const() && !(info->storage_class == ast::StorageClass::kNone)) {
-    diagnostics_.add_error("v-global01",
-                           "global constants shouldn't have a storage class",
-                           var->source());
+    AddError("v-global01", "global constants shouldn't have a storage class",
+             var->source());
     return false;
   }
 
@@ -670,9 +665,9 @@
   if (!ApplyStorageClassUsageToType(
           info->storage_class, const_cast<sem::Type*>(info->type->UnwrapRef()),
           var->source())) {
-    diagnostics_.add_note("while instantiating variable " +
-                              builder_->Symbols().NameFor(var->symbol()),
-                          var->source());
+    AddNote("while instantiating variable " +
+                builder_->Symbols().NameFor(var->symbol()),
+            var->source());
     return false;
   }
 
@@ -682,15 +677,13 @@
 bool Resolver::ValidateGlobalVariable(const VariableInfo* info) {
   auto duplicate_func = symbol_to_function_.find(info->declaration->symbol());
   if (duplicate_func != symbol_to_function_.end()) {
-    diagnostics_.add_error(
-        "v-2000",
-        "duplicate declaration '" +
-            builder_->Symbols().NameFor(info->declaration->symbol()) + "'",
-        info->declaration->source());
-    diagnostics_.add_note(
-        "'" + builder_->Symbols().NameFor(info->declaration->symbol()) +
-            "' first declared here:",
-        duplicate_func->second->declaration->source());
+    AddError("v-2000",
+             "duplicate declaration '" +
+                 builder_->Symbols().NameFor(info->declaration->symbol()) + "'",
+             info->declaration->source());
+    AddNote("'" + builder_->Symbols().NameFor(info->declaration->symbol()) +
+                "' first declared here:",
+            duplicate_func->second->declaration->source());
     return false;
   }
 
@@ -705,27 +698,23 @@
           uint32_t id = override_deco->value();
           auto itr = constant_ids_.find(id);
           if (itr != constant_ids_.end() && itr->second != info) {
-            diagnostics_.add_error("pipeline constant IDs must be unique",
-                                   deco->source());
-            diagnostics_.add_note("a pipeline constant with an ID of " +
-                                      std::to_string(id) +
-                                      " was previously declared "
-                                      "here:",
-                                  ast::GetDecoration<ast::OverrideDecoration>(
-                                      itr->second->declaration->decorations())
-                                      ->source());
+            AddError("pipeline constant IDs must be unique", deco->source());
+            AddNote("a pipeline constant with an ID of " + std::to_string(id) +
+                        " was previously declared "
+                        "here:",
+                    ast::GetDecoration<ast::OverrideDecoration>(
+                        itr->second->declaration->decorations())
+                        ->source());
             return false;
           }
           if (id > 65535) {
-            diagnostics_.add_error(
-                "pipeline constant IDs must be between 0 and 65535",
-                deco->source());
+            AddError("pipeline constant IDs must be between 0 and 65535",
+                     deco->source());
             return false;
           }
         }
       } else {
-        diagnostics_.add_error("decoration is not valid for constants",
-                               deco->source());
+        AddError("decoration is not valid for constants", deco->source());
         return false;
       }
     } else {
@@ -734,8 +723,7 @@
             deco->Is<ast::GroupDecoration>() ||
             deco->Is<ast::LocationDecoration>() ||
             deco->Is<ast::InternalDecoration>())) {
-        diagnostics_.add_error("decoration is not valid for variables",
-                               deco->source());
+        AddError("decoration is not valid for variables", deco->source());
         return false;
       }
     }
@@ -750,7 +738,7 @@
       // Each resource variable must be declared with both group and binding
       // attributes.
       if (!binding_point) {
-        diagnostics_.add_error(
+        AddError(
             "resource variables require [[group]] and [[binding]] "
             "decorations",
             info->declaration->source());
@@ -762,7 +750,7 @@
       if (binding_point.binding || binding_point.group) {
         // https://gpuweb.github.io/gpuweb/wgsl/#attribute-binding
         // Must only be applied to a resource variable
-        diagnostics_.add_error(
+        AddError(
             "non-resource variables must not have [[group]] or [[binding]] "
             "decorations",
             info->declaration->source());
@@ -775,7 +763,7 @@
   // storage storage class, must not be written.
   if (info->storage_class != ast::StorageClass::kStorage &&
       info->declaration->declared_access() != ast::Access::kUndefined) {
-    diagnostics_.add_error(
+    AddError(
         "variables declared not declared in the <storage> storage class must "
         "not declare an access control",
         info->declaration->source());
@@ -792,7 +780,7 @@
       auto* str = info->type->UnwrapRef()->As<sem::Struct>();
 
       if (!str) {
-        diagnostics_.add_error(
+        AddError(
             "variables declared in the <storage> storage class must be of a "
             "structure type",
             info->declaration->source());
@@ -800,13 +788,13 @@
       }
 
       if (!str->IsBlockDecorated()) {
-        diagnostics_.add_error(
+        AddError(
             "structure used as a storage buffer must be declared with the "
             "[[block]] decoration",
             str->Declaration()->source());
         if (info->declaration->source().range.begin.line) {
-          diagnostics_.add_note("structure used as storage buffer here",
-                                info->declaration->source());
+          AddNote("structure used as storage buffer here",
+                  info->declaration->source());
         }
         return false;
       }
@@ -819,7 +807,7 @@
       // attribute, satisfying the storage class constraints.
       auto* str = info->type->UnwrapRef()->As<sem::Struct>();
       if (!str) {
-        diagnostics_.add_error(
+        AddError(
             "variables declared in the <uniform> storage class must be of a "
             "structure type",
             info->declaration->source());
@@ -827,13 +815,13 @@
       }
 
       if (!str->IsBlockDecorated()) {
-        diagnostics_.add_error(
+        AddError(
             "structure used as a uniform buffer must be declared with the "
             "[[block]] decoration",
             str->Declaration()->source());
         if (info->declaration->source().range.begin.line) {
-          diagnostics_.add_note("structure used as uniform buffer here",
-                                info->declaration->source());
+          AddNote("structure used as uniform buffer here",
+                  info->declaration->source());
         }
         return false;
       }
@@ -841,12 +829,11 @@
       for (auto* member : str->Members()) {
         if (auto* arr = member->Type()->As<sem::Array>()) {
           if (arr->IsRuntimeSized()) {
-            diagnostics_.add_error(
+            AddError(
                 "structure containing a runtime sized array "
                 "cannot be used as a uniform buffer",
                 info->declaration->source());
-            diagnostics_.add_note("structure is declared here",
-                                  str->Declaration()->source());
+            AddNote("structure is declared here", str->Declaration()->source());
             return false;
           }
         }
@@ -866,33 +853,30 @@
   auto* storage_type = info->type->UnwrapRef();
 
   if (!var->is_const() && !IsStorable(storage_type)) {
-    diagnostics_.add_error(storage_type->FriendlyName(builder_->Symbols()) +
-                               " cannot be used as the type of a var",
-                           var->source());
+    AddError(storage_type->FriendlyName(builder_->Symbols()) +
+                 " cannot be used as the type of a var",
+             var->source());
     return false;
   }
 
   if (auto* r = storage_type->As<sem::Array>()) {
     if (r->IsRuntimeSized()) {
-      diagnostics_.add_error(
-          "v-0015",
-          "runtime arrays may only appear as the last member of a struct",
-          var->source());
+      AddError("v-0015",
+               "runtime arrays may only appear as the last member of a struct",
+               var->source());
       return false;
     }
   }
 
   if (auto* r = storage_type->As<sem::MultisampledTexture>()) {
     if (r->dim() != ast::TextureDimension::k2d) {
-      diagnostics_.add_error("only 2d multisampled textures are supported",
-                             var->source());
+      AddError("only 2d multisampled textures are supported", var->source());
       return false;
     }
 
     if (!r->type()->UnwrapRef()->is_numeric_scalar()) {
-      diagnostics_.add_error(
-          "texture_multisampled_2d<type>: type must be f32, i32 or u32",
-          var->source());
+      AddError("texture_multisampled_2d<type>: type must be f32, i32 or u32",
+               var->source());
       return false;
     }
   }
@@ -903,9 +887,9 @@
     // If the store type is a texture type or a sampler type, then the
     // variable declaration must not have a storage class decoration. The
     // storage class will always be handle.
-    diagnostics_.add_error("variables of type '" + info->type_name +
-                               "' must not have a storage class",
-                           var->source());
+    AddError("variables of type '" + info->type_name +
+                 "' must not have a storage class",
+             var->source());
     return false;
   }
 
@@ -916,15 +900,15 @@
     if (info->kind != VariableKind::kGlobal) {
       // Neither storage nor workgroup storage classes can be used in function
       // scopes.
-      diagnostics_.add_error("cannot declare an atomic var in a function scope",
-                             info->declaration->type()->source());
+      AddError("cannot declare an atomic var in a function scope",
+               info->declaration->type()->source());
       return false;
     }
     if (info->storage_class != ast::StorageClass::kWorkgroup) {
       // Storage buffers require a structure, so just check for workgroup
       // storage here.
-      diagnostics_.add_error("atomic var requires workgroup storage",
-                             info->declaration->type()->source());
+      AddError("atomic var requires workgroup storage",
+               info->declaration->type()->source());
       return false;
     }
   }
@@ -952,9 +936,8 @@
   switch (deco->value()) {
     case ast::Builtin::kPosition:
       if (!(type->is_float_vector() && type->As<sem::Vector>()->size() == 4)) {
-        diagnostics_.add_error(
-            "store type of " + deco_to_str(deco) + " must be 'vec4<f32>'",
-            deco->source());
+        AddError("store type of " + deco_to_str(deco) + " must be 'vec4<f32>'",
+                 deco->source());
         return false;
       }
       break;
@@ -963,25 +946,22 @@
     case ast::Builtin::kWorkgroupId:
       if (!(type->is_unsigned_integer_vector() &&
             type->As<sem::Vector>()->size() == 3)) {
-        diagnostics_.add_error(
-            "store type of " + deco_to_str(deco) + " must be 'vec3<u32>'",
-            deco->source());
+        AddError("store type of " + deco_to_str(deco) + " must be 'vec3<u32>'",
+                 deco->source());
         return false;
       }
       break;
     case ast::Builtin::kFragDepth:
       if (!type->Is<sem::F32>()) {
-        diagnostics_.add_error(
-            "store type of " + deco_to_str(deco) + " must be 'f32'",
-            deco->source());
+        AddError("store type of " + deco_to_str(deco) + " must be 'f32'",
+                 deco->source());
         return false;
       }
       break;
     case ast::Builtin::kFrontFacing:
       if (!type->Is<sem::Bool>()) {
-        diagnostics_.add_error(
-            "store type of " + deco_to_str(deco) + " must be 'bool'",
-            deco->source());
+        AddError("store type of " + deco_to_str(deco) + " must be 'bool'",
+                 deco->source());
         return false;
       }
       break;
@@ -991,9 +971,8 @@
     case ast::Builtin::kSampleMask:
     case ast::Builtin::kSampleIndex:
       if (!type->Is<sem::U32>()) {
-        diagnostics_.add_error(
-            "store type of " + deco_to_str(deco) + " must be 'u32'",
-            deco->source());
+        AddError("store type of " + deco_to_str(deco) + " must be 'u32'",
+                 deco->source());
         return false;
       }
       break;
@@ -1007,13 +986,12 @@
                                 const FunctionInfo* info) {
   auto func_it = symbol_to_function_.find(func->symbol());
   if (func_it != symbol_to_function_.end()) {
-    diagnostics_.add_error("v-0016",
-                           "duplicate function named '" +
-                               builder_->Symbols().NameFor(func->symbol()) +
-                               "'",
-                           func->source());
-    diagnostics_.add_note("first function declared here",
-                          func_it->second->declaration->source());
+    AddError("v-0016",
+             "duplicate function named '" +
+                 builder_->Symbols().NameFor(func->symbol()) + "'",
+             func->source());
+    AddNote("first function declared here",
+            func_it->second->declaration->source());
     return false;
   }
 
@@ -1021,14 +999,13 @@
   VariableInfo* var;
   if (variable_stack_.get(func->symbol(), &var, &is_global)) {
     if (is_global) {
-      diagnostics_.add_error("v-2000",
-                             "duplicate declaration '" +
-                                 builder_->Symbols().NameFor(func->symbol()) +
-                                 "'",
-                             func->source());
-      diagnostics_.add_note("'" + builder_->Symbols().NameFor(func->symbol()) +
-                                "' first declared here:",
-                            var->declaration->source());
+      AddError("v-2000",
+               "duplicate declaration '" +
+                   builder_->Symbols().NameFor(func->symbol()) + "'",
+               func->source());
+      AddNote("'" + builder_->Symbols().NameFor(func->symbol()) +
+                  "' first declared here:",
+              var->declaration->source());
       return false;
     }
   }
@@ -1041,14 +1018,13 @@
     } else if (deco->Is<ast::WorkgroupDecoration>()) {
       workgroup_deco_count++;
       if (func->pipeline_stage() != ast::PipelineStage::kCompute) {
-        diagnostics_.add_error(
+        AddError(
             "the workgroup_size attribute is only valid for compute stages",
             deco->source());
         return false;
       }
     } else if (!deco->Is<ast::InternalDecoration>()) {
-      diagnostics_.add_error("decoration is not valid for functions",
-                             deco->source());
+      AddError("decoration is not valid for functions", deco->source());
       return false;
     }
   }
@@ -1061,33 +1037,30 @@
 
   if (!info->return_type->Is<sem::Void>()) {
     if (!IsAtomicFreePlain(info->return_type)) {
-      diagnostics_.add_error(
-          "function return type must be an atomic-free plain type",
-          func->return_type()->source());
+      AddError("function return type must be an atomic-free plain type",
+               func->return_type()->source());
       return false;
     }
 
     if (func->body()) {
       if (!func->get_last_statement() ||
           !func->get_last_statement()->Is<ast::ReturnStatement>()) {
-        diagnostics_.add_error(
-            "v-0002", "non-void function must end with a return statement",
-            func->source());
+        AddError("v-0002", "non-void function must end with a return statement",
+                 func->source());
         return false;
       }
     } else if (!IsValidationDisabled(
                    func->decorations(),
                    ast::DisabledValidation::kFunctionHasNoBody)) {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Resolver, diagnostics_)
           << "Function " << builder_->Symbols().NameFor(func->symbol())
           << " has no body";
     }
 
     for (auto* deco : func->return_type_decorations()) {
       if (!deco->IsAnyOf<ast::BuiltinDecoration, ast::LocationDecoration>()) {
-        diagnostics_.add_error(
-            "decoration is not valid for function return types",
-            deco->source());
+        AddError("decoration is not valid for function return types",
+                 deco->source());
         return false;
       }
     }
@@ -1126,9 +1099,8 @@
         for (auto* deco : decos) {
           if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
             if (pipeline_io_attribute) {
-              diagnostics_.add_error("multiple entry point IO attributes",
-                                     deco->source());
-              diagnostics_.add_note(
+              AddError("multiple entry point IO attributes", deco->source());
+              AddNote(
                   "previously consumed " + deco_to_str(pipeline_io_attribute),
                   pipeline_io_attribute->source());
               return false;
@@ -1136,7 +1108,7 @@
             pipeline_io_attribute = deco;
 
             if (builtins.count(builtin->value())) {
-              diagnostics_.add_error(
+              AddError(
                   deco_to_str(builtin) +
                       " attribute appears multiple times as pipeline " +
                       (param_or_ret == ParamOrRetType::kParameter ? "input"
@@ -1148,9 +1120,8 @@
 
           } else if (auto* location = deco->As<ast::LocationDecoration>()) {
             if (pipeline_io_attribute) {
-              diagnostics_.add_error("multiple entry point IO attributes",
-                                     deco->source());
-              diagnostics_.add_note(
+              AddError("multiple entry point IO attributes", deco->source());
+              AddNote(
                   "previously consumed " + deco_to_str(pipeline_io_attribute),
                   pipeline_io_attribute->source());
               return false;
@@ -1158,7 +1129,7 @@
             pipeline_io_attribute = deco;
 
             if (locations.count(location->value())) {
-              diagnostics_.add_error(
+              AddError(
                   deco_to_str(location) +
                       " attribute appears multiple times as pipeline " +
                       (param_or_ret == ParamOrRetType::kParameter ? "input"
@@ -1173,7 +1144,7 @@
         // Check that we saw a pipeline IO attribute iff we need one.
         if (ty->Is<sem::Struct>()) {
           if (pipeline_io_attribute) {
-            diagnostics_.add_error(
+            AddError(
                 "entry point IO attributes must not be used on structure " +
                     std::string(param_or_ret == ParamOrRetType::kParameter
                                     ? "parameters"
@@ -1190,7 +1161,7 @@
                           ? " on parameter"
                           : " on return type");
             }
-            diagnostics_.add_error(err, source);
+            AddError(err, source);
             return false;
           }
 
@@ -1199,7 +1170,7 @@
           // Testing for being a struct is handled by the if portion above.
           if (!pipeline_io_attribute->Is<ast::BuiltinDecoration>()) {
             if (!ty->is_numeric_scalar_or_vector()) {
-              diagnostics_.add_error(
+              AddError(
                   "User defined entry point IO types must be a numeric scalar, "
                   "a numeric vector, or a structure",
                   source);
@@ -1226,24 +1197,21 @@
       // invalid member types.
       for (auto* member : str->Members()) {
         if (member->Type()->Is<sem::Struct>()) {
-          diagnostics_.add_error(
-              "entry point IO types cannot contain nested structures",
-              member->Declaration()->source());
-          diagnostics_.add_note("while analysing entry point " +
-                                    builder_->Symbols().NameFor(func->symbol()),
-                                func->source());
+          AddError("entry point IO types cannot contain nested structures",
+                   member->Declaration()->source());
+          AddNote("while analysing entry point " +
+                      builder_->Symbols().NameFor(func->symbol()),
+                  func->source());
           return false;
         }
 
         if (auto* arr = member->Type()->As<sem::Array>()) {
           if (arr->IsRuntimeSized()) {
-            diagnostics_.add_error(
-                "entry point IO types cannot contain runtime sized arrays",
-                member->Declaration()->source());
-            diagnostics_.add_note(
-                "while analysing entry point " +
-                    builder_->Symbols().NameFor(func->symbol()),
-                func->source());
+            AddError("entry point IO types cannot contain runtime sized arrays",
+                     member->Declaration()->source());
+            AddNote("while analysing entry point " +
+                        builder_->Symbols().NameFor(func->symbol()),
+                    func->source());
             return false;
           }
         }
@@ -1251,9 +1219,9 @@
         if (!validate_entry_point_decorations_inner(
                 member->Declaration()->decorations(), member->Type(),
                 member->Declaration()->source(), param_or_ret, true)) {
-          diagnostics_.add_note("while analysing entry point " +
-                                    builder_->Symbols().NameFor(func->symbol()),
-                                func->source());
+          AddNote("while analysing entry point " +
+                      builder_->Symbols().NameFor(func->symbol()),
+                  func->source());
           return false;
         }
       }
@@ -1298,7 +1266,7 @@
       }
     }
     if (!found) {
-      diagnostics_.add_error(
+      AddError(
           "a vertex shader must include the 'position' builtin in its return "
           "type",
           func->source());
@@ -1327,14 +1295,14 @@
       // the same group and binding values, when considered as a pair of
       // values.
       auto func_name = builder_->Symbols().NameFor(info->declaration->symbol());
-      diagnostics_.add_error("entry point '" + func_name +
-                                 "' references multiple variables that use the "
-                                 "same resource binding [[group(" +
-                                 std::to_string(bp.group) + "), binding(" +
-                                 std::to_string(bp.binding) + ")]]",
-                             var_info->declaration->source());
-      diagnostics_.add_note("first resource binding usage declared here",
-                            res.first->second->source());
+      AddError("entry point '" + func_name +
+                   "' references multiple variables that use the "
+                   "same resource binding [[group(" +
+                   std::to_string(bp.group) + "), binding(" +
+                   std::to_string(bp.binding) + ")]]",
+               var_info->declaration->source());
+      AddNote("first resource binding usage declared here",
+              res.first->second->source());
       return false;
     }
   }
@@ -1372,9 +1340,9 @@
 
     if (!ApplyStorageClassUsageToType(param->declared_storage_class(),
                                       param_info->type, param->source())) {
-      diagnostics_.add_note("while instantiating parameter " +
-                                builder_->Symbols().NameFor(param->symbol()),
-                            param->source());
+      AddNote("while instantiating parameter " +
+                  builder_->Symbols().NameFor(param->symbol()),
+              param->source());
       return false;
     }
 
@@ -1410,9 +1378,9 @@
   if (auto* str = info->return_type->As<sem::Struct>()) {
     if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, str,
                                       func->source())) {
-      diagnostics_.add_note("while instantiating return type for " +
-                                builder_->Symbols().NameFor(func->symbol()),
-                            func->source());
+      AddNote("while instantiating return type for " +
+                  builder_->Symbols().NameFor(func->symbol()),
+              func->source());
       return false;
     }
 
@@ -1434,7 +1402,7 @@
   if (func->body()) {
     Mark(func->body());
     if (current_statement_) {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Resolver, diagnostics_)
           << "Resolver::Function() called with a current statement";
       return false;
     }
@@ -1492,7 +1460,7 @@
         VariableInfo* var;
         if (!variable_stack_.get(ident->symbol(), &var) ||
             !(var->declaration->is_const() && var->type->Is<sem::I32>())) {
-          diagnostics_.add_error(
+          AddError(
               "workgroup_size parameter must be a literal i32 or an i32 "
               "module-scope constant",
               values[i]->source());
@@ -1522,7 +1490,7 @@
         Mark(scalar->literal());
 
         if (!scalar->literal()->Is<ast::IntLiteral>()) {
-          diagnostics_.add_error(
+          AddError(
               "workgroup_size parameter must be a literal i32 or an i32 "
               "module-scope constant",
               values[i]->source());
@@ -1536,9 +1504,8 @@
 
       // Validate and set the default value for this dimension.
       if (value < 1) {
-        diagnostics_.add_error(
-            "workgroup_size parameter must be a positive i32 value",
-            values[i]->source());
+        AddError("workgroup_size parameter must be a positive i32 value",
+                 values[i]->source());
         return false;
       }
       info->workgroup_size[i].value = value;
@@ -1579,7 +1546,7 @@
     if (stmt->IsAnyOf<ast::ReturnStatement, ast::BreakStatement,
                       ast::ContinueStatement>()) {
       if (stmt != stmts.back()) {
-        diagnostics_.add_error("code is unreachable", (*next_stmt)->source());
+        AddError("code is unreachable", (*next_stmt)->source());
         return false;
       }
     }
@@ -1600,7 +1567,7 @@
   TINT_SCOPED_ASSIGNMENT(current_statement_, sem_statement);
 
   if (stmt->Is<ast::ElseStatement>()) {
-    TINT_ICE(diagnostics_)
+    TINT_ICE(Resolver, diagnostics_)
         << "Resolver::Statement() encountered an Else statement. Else "
            "statements are embedded in If statements, so should never be "
            "encountered as top-level statements";
@@ -1616,8 +1583,8 @@
   if (stmt->Is<ast::BreakStatement>()) {
     if (!current_block_->FindFirstParent<sem::LoopBlockStatement>() &&
         !current_block_->FindFirstParent<sem::SwitchCaseBlockStatement>()) {
-      diagnostics_.add_error("break statement must be in a loop or switch case",
-                             stmt->source());
+      AddError("break statement must be in a loop or switch case",
+               stmt->source());
       return false;
     }
     return true;
@@ -1644,8 +1611,7 @@
             ->SetFirstContinue(loop_block->Decls().size());
       }
     } else {
-      diagnostics_.add_error("continue statement must be in a loop",
-                             stmt->source());
+      AddError("continue statement must be in a loop", stmt->source());
       return false;
     }
 
@@ -1673,7 +1639,7 @@
     return VariableDeclStatement(v);
   }
 
-  diagnostics_.add_error(
+  AddError(
       "unknown statement type for type determination: " + builder_->str(stmt),
       stmt->source());
   return false;
@@ -1700,9 +1666,9 @@
 
   auto* cond_type = TypeOf(stmt->condition())->UnwrapRef();
   if (!cond_type->Is<sem::Bool>()) {
-    diagnostics_.add_error("if statement condition must be bool, got " +
-                               cond_type->FriendlyName(builder_->Symbols()),
-                           stmt->condition()->source());
+    AddError("if statement condition must be bool, got " +
+                 cond_type->FriendlyName(builder_->Symbols()),
+             stmt->condition()->source());
     return false;
   }
 
@@ -1732,10 +1698,9 @@
 
       auto* else_cond_type = TypeOf(cond)->UnwrapRef();
       if (!else_cond_type->Is<sem::Bool>()) {
-        diagnostics_.add_error(
-            "else statement condition must be bool, got " +
-                else_cond_type->FriendlyName(builder_->Symbols()),
-            cond->source());
+        AddError("else statement condition must be bool, got " +
+                     else_cond_type->FriendlyName(builder_->Symbols()),
+                 cond->source());
         return false;
       }
     }
@@ -1820,8 +1785,7 @@
   } else if (auto* unary = expr->As<ast::UnaryOpExpression>()) {
     ok = UnaryOp(unary);
   } else {
-    diagnostics_.add_error("unknown expression for type determination",
-                           expr->source());
+    AddError("unknown expression for type determination", expr->source());
   }
 
   if (!ok) {
@@ -1830,8 +1794,8 @@
 
   auto* ty = TypeOf(expr);
   if (ty->Is<sem::Atomic>()) {
-    diagnostics_.add_error("an expression must not evaluate to an atomic type",
-                           expr->source());
+    AddError("an expression must not evaluate to an atomic type",
+             expr->source());
     return false;
   }
 
@@ -1859,16 +1823,16 @@
   } else if (auto* mat = parent_type->As<sem::Matrix>()) {
     ret = builder_->create<sem::Vector>(mat->type(), mat->rows());
   } else {
-    diagnostics_.add_error("invalid parent type (" + parent_type->type_name() +
-                               ") in array accessor",
-                           expr->source());
+    AddError("invalid parent type (" + parent_type->type_name() +
+                 ") in array accessor",
+             expr->source());
     return false;
   }
 
   if (!TypeOf(idx)->UnwrapRef()->IsAnyOf<sem::I32, sem::U32>()) {
-    diagnostics_.add_error("index must be of type 'i32' or 'u32', found: '" +
-                               TypeNameOf(idx) + "'",
-                           idx->source());
+    AddError("index must be of type 'i32' or 'u32', found: '" +
+                 TypeNameOf(idx) + "'",
+             idx->source());
     return false;
   }
 
@@ -1878,8 +1842,8 @@
       // https://github.com/gpuweb/gpuweb/issues/1272
       auto* scalar = idx->As<ast::ScalarConstructorExpression>();
       if (!scalar || !scalar->literal()->As<ast::IntLiteral>()) {
-        diagnostics_.add_error(
-            "index must be signed or unsigned integer literal", idx->source());
+        AddError("index must be signed or unsigned integer literal",
+                 idx->source());
         return false;
       }
     }
@@ -1905,7 +1869,7 @@
     return false;
   }
   if (ty->Is<sem::Pointer>()) {
-    diagnostics_.add_error("cannot cast to a pointer", expr->source());
+    AddError("cannot cast to a pointer", expr->source());
     return false;
   }
   SetType(expr, ty, expr->type()->FriendlyName(builder_->Symbols()));
@@ -1948,8 +1912,9 @@
     if (auto* intrinsic = target->As<sem::Intrinsic>()) {
       return_type = intrinsic->ReturnType();
     } else {
-      TINT_ICE(diagnostics_) << "call target was not an intrinsic, but a "
-                             << intrinsic->TypeInfo().name;
+      TINT_ICE(Resolver, diagnostics_)
+          << "call target was not an intrinsic, but a "
+          << intrinsic->TypeInfo().name;
     }
   }
 
@@ -1960,7 +1925,7 @@
     // that value must be consumed either through assignment, evaluation in
     // another expression or through use of the ignore built-in function (see
     // § 16.13 Value-steering functions).
-    diagnostics_.add_error(
+    AddError(
         "result of called function was not used. If this was intentional wrap "
         "the function call in ignore()",
         stmt->source());
@@ -1984,7 +1949,7 @@
   }
 
   if (result->IsDeprecated()) {
-    diagnostics_.add_warning("use of deprecated intrinsic", call->source());
+    AddWarning("use of deprecated intrinsic", call->source());
   }
 
   builder_->Sem().Add(
@@ -2005,13 +1970,13 @@
   if (callee_func_it == symbol_to_function_.end()) {
     if (current_function_ &&
         current_function_->declaration->symbol() == ident->symbol()) {
-      diagnostics_.add_error("v-0004",
-                             "recursion is not permitted. '" + name +
-                                 "' attempted to call itself.",
-                             call->source());
+      AddError("v-0004",
+               "recursion is not permitted. '" + name +
+                   "' attempted to call itself.",
+               call->source());
     } else {
-      diagnostics_.add_error("v-0006: unable to find called function: " + name,
-                             call->source());
+      AddError("v-0006: unable to find called function: " + name,
+               call->source());
     }
     return false;
   }
@@ -2037,12 +2002,11 @@
   // Validate number of arguments match number of parameters
   if (call->params().size() != callee_func->parameters.size()) {
     bool more = call->params().size() > callee_func->parameters.size();
-    diagnostics_.add_error(
-        "too " + (more ? std::string("many") : std::string("few")) +
-            " arguments in call to '" + name + "', expected " +
-            std::to_string(callee_func->parameters.size()) + ", got " +
-            std::to_string(call->params().size()),
-        call->source());
+    AddError("too " + (more ? std::string("many") : std::string("few")) +
+                 " arguments in call to '" + name + "', expected " +
+                 std::to_string(callee_func->parameters.size()) + ", got " +
+                 std::to_string(call->params().size()),
+             call->source());
     return false;
   }
 
@@ -2053,12 +2017,11 @@
     auto* arg_type = TypeOf(arg_expr)->UnwrapRef();
 
     if (param->type != arg_type) {
-      diagnostics_.add_error(
-          "type mismatch for argument " + std::to_string(i + 1) +
-              " in call to '" + name + "', expected '" +
-              param->type->FriendlyName(builder_->Symbols()) + "', got '" +
-              arg_type->FriendlyName(builder_->Symbols()) + "'",
-          arg_expr->source());
+      AddError("type mismatch for argument " + std::to_string(i + 1) +
+                   " in call to '" + name + "', expected '" +
+                   param->type->FriendlyName(builder_->Symbols()) + "', got '" +
+                   arg_type->FriendlyName(builder_->Symbols()) + "'",
+               arg_expr->source());
       return false;
     }
   }
@@ -2090,7 +2053,7 @@
     // obey the constructor type rules laid out in
     // https://gpuweb.github.io/gpuweb/wgsl.html#type-constructor-expr.
     if (type->Is<sem::Pointer>()) {
-      diagnostics_.add_error("cannot cast to a pointer", expr->source());
+      AddError("cannot cast to a pointer", expr->source());
       return false;
     }
     if (auto* vec_type = type->As<sem::Vector>()) {
@@ -2113,7 +2076,8 @@
     }
     SetType(expr, type);
   } else {
-    TINT_ICE(diagnostics_) << "unexpected constructor expression type";
+    TINT_ICE(Resolver, diagnostics_)
+        << "unexpected constructor expression type";
   }
   return true;
 }
@@ -2126,7 +2090,7 @@
   for (auto* value : values) {
     auto* value_type = TypeOf(value)->UnwrapRef();
     if (value_type != elem_type) {
-      diagnostics_.add_error(
+      AddError(
           "type in array constructor does not match array type: "
           "expected '" +
               elem_type->FriendlyName(builder_->Symbols()) + "', found '" +
@@ -2137,22 +2101,20 @@
   }
 
   if (array_type->IsRuntimeSized()) {
-    diagnostics_.add_error("cannot init a runtime-sized array", ctor->source());
+    AddError("cannot init a runtime-sized array", ctor->source());
     return false;
   } else if (!values.empty() && (values.size() != array_type->Count())) {
     std::string fm = values.size() < array_type->Count() ? "few" : "many";
-    diagnostics_.add_error("array constructor has too " + fm +
-                               " elements: expected " +
-                               std::to_string(array_type->Count()) +
-                               ", found " + std::to_string(values.size()),
-                           ctor->source());
+    AddError("array constructor has too " + fm + " elements: expected " +
+                 std::to_string(array_type->Count()) + ", found " +
+                 std::to_string(values.size()),
+             ctor->source());
     return false;
   } else if (values.size() > array_type->Count()) {
-    diagnostics_.add_error(
-        "array constructor has too many elements: expected " +
-            std::to_string(array_type->Count()) + ", found " +
-            std::to_string(values.size()),
-        ctor->source());
+    AddError("array constructor has too many elements: expected " +
+                 std::to_string(array_type->Count()) + ", found " +
+                 std::to_string(values.size()),
+             ctor->source());
     return false;
   }
   return true;
@@ -2168,7 +2130,7 @@
     auto* value_type = TypeOf(value)->UnwrapRef();
     if (value_type->is_scalar()) {
       if (elem_type != value_type) {
-        diagnostics_.add_error(
+        AddError(
             "type in vector constructor does not match vector type: "
             "expected '" +
                 elem_type->FriendlyName(builder_->Symbols()) + "', found '" +
@@ -2188,7 +2150,7 @@
       // https://gpuweb.github.io/gpuweb/wgsl.html#conversion-expr).
       if (elem_type != value_elem_type &&
           (values.size() > 1u || value_vec->is_bool_vector())) {
-        diagnostics_.add_error(
+        AddError(
             "type in vector constructor does not match vector type: "
             "expected '" +
                 elem_type->FriendlyName(builder_->Symbols()) + "', found '" +
@@ -2200,10 +2162,9 @@
       value_cardinality_sum += value_vec->size();
     } else {
       // A vector constructor can only accept vectors and scalars.
-      diagnostics_.add_error(
-          "expected vector or scalar type in vector constructor; found: " +
-              value_type->FriendlyName(builder_->Symbols()),
-          value->source());
+      AddError("expected vector or scalar type in vector constructor; found: " +
+                   value_type->FriendlyName(builder_->Symbols()),
+               value->source());
       return false;
     }
   }
@@ -2213,15 +2174,14 @@
   // of all constructor arguments must add up to the vector cardinality.
   if (value_cardinality_sum > 1 && value_cardinality_sum != vec_type->size()) {
     if (values.empty()) {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Resolver, diagnostics_)
           << "constructor arguments expected to be non-empty!";
     }
     const Source& values_start = values[0]->source();
     const Source& values_end = values[values.size() - 1]->source();
-    diagnostics_.add_error(
-        "attempted to construct '" + TypeNameOf(ctor) + "' with " +
-            std::to_string(value_cardinality_sum) + " component(s)",
-        Source::Combine(values_start, values_end));
+    AddError("attempted to construct '" + TypeNameOf(ctor) + "' with " +
+                 std::to_string(value_cardinality_sum) + " component(s)",
+             Source::Combine(values_start, values_end));
     return false;
   }
   return true;
@@ -2230,7 +2190,7 @@
 bool Resolver::ValidateMatrix(const sem::Matrix* matrix_type,
                               const Source& source) {
   if (!matrix_type->is_float_matrix()) {
-    diagnostics_.add_error("matrix element type must be 'f32'", source);
+    AddError("matrix element type must be 'f32'", source);
     return false;
   }
   return true;
@@ -2253,12 +2213,11 @@
   if (matrix_type->columns() != values.size()) {
     const Source& values_start = values[0]->source();
     const Source& values_end = values[values.size() - 1]->source();
-    diagnostics_.add_error(
-        "expected " + std::to_string(matrix_type->columns()) + " '" +
-            VectorPretty(matrix_type->rows(), elem_type) + "' arguments in '" +
-            TypeNameOf(ctor) + "' constructor, found " +
-            std::to_string(values.size()),
-        Source::Combine(values_start, values_end));
+    AddError("expected " + std::to_string(matrix_type->columns()) + " '" +
+                 VectorPretty(matrix_type->rows(), elem_type) +
+                 "' arguments in '" + TypeNameOf(ctor) +
+                 "' constructor, found " + std::to_string(values.size()),
+             Source::Combine(values_start, values_end));
     return false;
   }
 
@@ -2268,12 +2227,11 @@
 
     if (!value_vec || value_vec->size() != matrix_type->rows() ||
         elem_type != value_vec->type()) {
-      diagnostics_.add_error("expected argument type '" +
-                                 VectorPretty(matrix_type->rows(), elem_type) +
-                                 "' in '" + TypeNameOf(ctor) +
-                                 "' constructor, found '" + TypeNameOf(value) +
-                                 "'",
-                             value->source());
+      AddError("expected argument type '" +
+                   VectorPretty(matrix_type->rows(), elem_type) + "' in '" +
+                   TypeNameOf(ctor) + "' constructor, found '" +
+                   TypeNameOf(value) + "'",
+               value->source());
       return false;
     }
   }
@@ -2288,9 +2246,9 @@
     return true;
   }
   if (ctor->values().size() > 1) {
-    diagnostics_.add_error("expected zero or one value in constructor, got " +
-                               std::to_string(ctor->values().size()),
-                           ctor->source());
+    AddError("expected zero or one value in constructor, got " +
+                 std::to_string(ctor->values().size()),
+             ctor->source());
     return false;
   }
 
@@ -2309,10 +2267,9 @@
       (type->Is<U32>() && value_type->IsAnyOf<I32, U32, F32>()) ||
       (type->Is<F32>() && value_type->IsAnyOf<I32, U32, F32>());
   if (!is_valid) {
-    diagnostics_.add_error("cannot construct '" + TypeNameOf(ctor) +
-                               "' with a value of type '" + TypeNameOf(value) +
-                               "'",
-                           ctor->source());
+    AddError("cannot construct '" + TypeNameOf(ctor) +
+                 "' with a value of type '" + TypeNameOf(value) + "'",
+             ctor->source());
 
     return false;
   }
@@ -2349,11 +2306,10 @@
             auto var_decl_index =
                 static_cast<size_t>(std::distance(decls.begin(), iter));
             if (var_decl_index >= loop_block->FirstContinue()) {
-              diagnostics_.add_error(
-                  "continue statement bypasses declaration of '" +
-                      builder_->Symbols().NameFor(symbol) +
-                      "' in continuing block",
-                  expr->source());
+              AddError("continue statement bypasses declaration of '" +
+                           builder_->Symbols().NameFor(symbol) +
+                           "' in continuing block",
+                       expr->source());
               return false;
             }
           }
@@ -2366,21 +2322,18 @@
 
   auto iter = symbol_to_function_.find(symbol);
   if (iter != symbol_to_function_.end()) {
-    diagnostics_.add_error("missing '(' for function call",
-                           expr->source().End());
+    AddError("missing '(' for function call", expr->source().End());
     return false;
   }
 
   std::string name = builder_->Symbols().NameFor(symbol);
   if (sem::ParseIntrinsicType(name) != IntrinsicType::kNone) {
-    diagnostics_.add_error("missing '(' for intrinsic call",
-                           expr->source().End());
+    AddError("missing '(' for intrinsic call", expr->source().End());
     return false;
   }
 
-  diagnostics_.add_error(
-      "v-0006: identifier must be declared before use: " + name,
-      expr->source());
+  AddError("v-0006: identifier must be declared before use: " + name,
+           expr->source());
   return false;
 }
 
@@ -2410,7 +2363,7 @@
     }
 
     if (ret == nullptr) {
-      diagnostics_.add_error(
+      AddError(
           "struct member " + builder_->Symbols().NameFor(symbol) + " not found",
           expr->source());
       return false;
@@ -2449,22 +2402,19 @@
           swizzle.emplace_back(3);
           break;
         default:
-          diagnostics_.add_error(
-              "invalid vector swizzle character",
-              expr->member()->source().Begin() + swizzle.size());
+          AddError("invalid vector swizzle character",
+                   expr->member()->source().Begin() + swizzle.size());
           return false;
       }
 
       if (swizzle.back() >= vec->size()) {
-        diagnostics_.add_error("invalid vector swizzle member",
-                               expr->member()->source());
+        AddError("invalid vector swizzle member", expr->member()->source());
         return false;
       }
     }
 
     if (size < 1 || size > 4) {
-      diagnostics_.add_error("invalid vector swizzle size",
-                             expr->member()->source());
+      AddError("invalid vector swizzle size", expr->member()->source());
       return false;
     }
 
@@ -2477,9 +2427,8 @@
     };
     if (!std::all_of(s.begin(), s.end(), is_rgba) &&
         !std::all_of(s.begin(), s.end(), is_xyzw)) {
-      diagnostics_.add_error(
-          "invalid mixing of vector swizzle characters rgba with xyzw",
-          expr->member()->source());
+      AddError("invalid mixing of vector swizzle characters rgba with xyzw",
+               expr->member()->source());
       return false;
     }
 
@@ -2501,10 +2450,9 @@
         expr, builder_->create<sem::Swizzle>(expr, ret, current_statement_,
                                              std::move(swizzle)));
   } else {
-    diagnostics_.add_error(
-        "invalid use of member accessor on a non-vector/non-struct " +
-            TypeNameOf(expr->structure()),
-        expr->source());
+    AddError("invalid use of member accessor on a non-vector/non-struct " +
+                 TypeNameOf(expr->structure()),
+             expr->source());
     return false;
   }
 
@@ -2717,12 +2665,11 @@
     }
   }
 
-  diagnostics_.add_error(
-      "Binary expression operand types are invalid for this operation: " +
-          lhs_type->FriendlyName(builder_->Symbols()) + " " +
-          FriendlyName(expr->op()) + " " +
-          rhs_type->FriendlyName(builder_->Symbols()),
-      expr->source());
+  AddError("Binary expression operand types are invalid for this operation: " +
+               lhs_type->FriendlyName(builder_->Symbols()) + " " +
+               FriendlyName(expr->op()) + " " +
+               rhs_type->FriendlyName(builder_->Symbols()),
+           expr->source());
   return false;
 }
 
@@ -2756,8 +2703,8 @@
         type = builder_->create<sem::Pointer>(
             ref->StoreType(), ref->StorageClass(), ref->Access());
       } else {
-        diagnostics_.add_error("cannot take the address of expression",
-                               unary->expr()->source());
+        AddError("cannot take the address of expression",
+                 unary->expr()->source());
         return false;
       }
       break;
@@ -2767,9 +2714,9 @@
         type = builder_->create<sem::Reference>(
             ptr->StoreType(), ptr->StorageClass(), ptr->Access());
       } else {
-        diagnostics_.add_error("cannot dereference expression of type '" +
-                                   TypeNameOf(unary->expr()) + "'",
-                               unary->expr()->source());
+        AddError("cannot dereference expression of type '" +
+                     TypeNameOf(unary->expr()) + "'",
+                 unary->expr()->source());
         return false;
       }
       break;
@@ -2786,10 +2733,10 @@
   bool is_global = false;
   if (variable_stack_.get(var->symbol(), nullptr, &is_global)) {
     const char* error_code = is_global ? "v-0013" : "v-0014";
-    diagnostics_.add_error(error_code,
-                           "redeclared identifier '" +
-                               builder_->Symbols().NameFor(var->symbol()) + "'",
-                           var->source());
+    AddError(error_code,
+             "redeclared identifier '" +
+                 builder_->Symbols().NameFor(var->symbol()) + "'",
+             var->source());
     return false;
   }
 
@@ -2816,9 +2763,8 @@
             var->decorations(),
             ast::DisabledValidation::kFunctionVarStorageClass)) {
       if (info->storage_class != ast::StorageClass::kNone) {
-        diagnostics_.add_error(
-            "function variable has a non-function storage class",
-            stmt->source());
+        AddError("function variable has a non-function storage class",
+                 stmt->source());
         return false;
       }
       info->storage_class = ast::StorageClass::kFunction;
@@ -2827,9 +2773,9 @@
 
   if (!ApplyStorageClassUsageToType(info->storage_class, info->type,
                                     var->source())) {
-    diagnostics_.add_note("while instantiating variable " +
-                              builder_->Symbols().NameFor(var->symbol()),
-                          var->source());
+    AddNote("while instantiating variable " +
+                builder_->Symbols().NameFor(var->symbol()),
+            var->source());
     return false;
   }
 
@@ -2843,7 +2789,7 @@
   } else if (auto* str = named_type->As<ast::Struct>()) {
     result = Structure(str);
   } else {
-    TINT_UNREACHABLE(diagnostics_) << "Unhandled TypeDecl";
+    TINT_UNREACHABLE(Resolver, diagnostics_) << "Unhandled TypeDecl";
   }
 
   if (!result) {
@@ -2864,14 +2810,15 @@
 bool Resolver::ValidateTypeDecl(const ast::TypeDecl* named_type) const {
   auto iter = named_type_info_.find(named_type->name());
   if (iter == named_type_info_.end()) {
-    TINT_ICE(diagnostics_) << "ValidateTypeDecl called() before TypeDecl()";
+    TINT_ICE(Resolver, diagnostics_)
+        << "ValidateTypeDecl called() before TypeDecl()";
   }
   if (iter->second.ast != named_type) {
-    diagnostics_.add_error("type with the name '" +
-                               builder_->Symbols().NameFor(named_type->name()) +
-                               "' was already declared",
-                           named_type->source());
-    diagnostics_.add_note("first declared here", iter->second.ast->source());
+    AddError("type with the name '" +
+                 builder_->Symbols().NameFor(named_type->name()) +
+                 "' was already declared",
+             named_type->source());
+    AddNote("first declared here", iter->second.ast->source());
     return false;
   }
   return true;
@@ -2906,7 +2853,7 @@
   if (lit->Is<ast::BoolLiteral>()) {
     return builder_->create<sem::Bool>();
   }
-  TINT_UNREACHABLE(diagnostics_)
+  TINT_UNREACHABLE(Resolver, diagnostics_)
       << "Unhandled literal type: " << lit->TypeInfo().name;
   return nullptr;
 }
@@ -2919,7 +2866,8 @@
                        const sem::Type* type,
                        const std::string& type_name) {
   if (expr_info_.count(expr)) {
-    TINT_ICE(diagnostics_) << "SetType() called twice for the same expression";
+    TINT_ICE(Resolver, diagnostics_)
+        << "SetType() called twice for the same expression";
   }
   expr_info_.emplace(expr, ExpressionInfo{type, type_name, current_statement_});
 }
@@ -2938,28 +2886,27 @@
             if (it != expr_info_.end()) {
               if (func->declaration->symbol() ==
                   it->second.statement->Function()->symbol()) {
-                diagnostics_.add_error("workgroup memory cannot be used by " +
-                                           stage_name.str() + " pipeline stage",
-                                       user->source());
+                AddError("workgroup memory cannot be used by " +
+                             stage_name.str() + " pipeline stage",
+                         user->source());
                 break;
               }
             }
           }
-          diagnostics_.add_note("variable is declared here",
-                                var->declaration->source());
+          AddNote("variable is declared here", var->declaration->source());
           if (func != entry_point) {
             TraverseCallChain(entry_point, func, [&](FunctionInfo* f) {
-              diagnostics_.add_note(
+              AddNote(
                   "called by function '" +
                       builder_->Symbols().NameFor(f->declaration->symbol()) +
                       "'",
                   f->declaration->source());
             });
-            diagnostics_.add_note("called by entry point '" +
-                                      builder_->Symbols().NameFor(
-                                          entry_point->declaration->symbol()) +
-                                      "'",
-                                  entry_point->declaration->source());
+            AddNote("called by entry point '" +
+                        builder_->Symbols().NameFor(
+                            entry_point->declaration->symbol()) +
+                        "'",
+                    entry_point->declaration->source());
           }
           return false;
         }
@@ -2986,19 +2933,19 @@
       if (!call.intrinsic->SupportedStages().Contains(stage)) {
         std::stringstream err;
         err << "built-in cannot be used by " << stage << " pipeline stage";
-        diagnostics_.add_error(err.str(), call.call->source());
+        AddError(err.str(), call.call->source());
         if (func != entry_point) {
           TraverseCallChain(entry_point, func, [&](FunctionInfo* f) {
-            diagnostics_.add_note(
-                "called by function '" +
-                    builder_->Symbols().NameFor(f->declaration->symbol()) + "'",
-                f->declaration->source());
+            AddNote("called by function '" +
+                        builder_->Symbols().NameFor(f->declaration->symbol()) +
+                        "'",
+                    f->declaration->source());
           });
-          diagnostics_.add_note("called by entry point '" +
-                                    builder_->Symbols().NameFor(
-                                        entry_point->declaration->symbol()) +
-                                    "'",
-                                entry_point->declaration->source());
+          AddNote("called by entry point '" +
+                      builder_->Symbols().NameFor(
+                          entry_point->declaration->symbol()) +
+                      "'",
+                  entry_point->declaration->source());
         }
         return false;
       }
@@ -3034,7 +2981,7 @@
       return;
     }
   }
-  TINT_ICE(diagnostics_)
+  TINT_ICE(Resolver, diagnostics_)
       << "TraverseCallChain() 'from' does not transitively call 'to'";
 }
 
@@ -3072,7 +3019,7 @@
         constant_id = next_constant_id;
         while (constant_ids_.count(constant_id)) {
           if (constant_id == UINT16_MAX) {
-            TINT_ICE(builder_->Diagnostics())
+            TINT_ICE(Resolver, builder_->Diagnostics())
                 << "no more pipeline constant IDs available";
             return;
           }
@@ -3102,8 +3049,8 @@
       } else {
         auto* sem_user = sem_expr->As<sem::VariableUser>();
         if (!sem_user) {
-          TINT_ICE(diagnostics_) << "expected sem::VariableUser, got "
-                                 << sem_expr->TypeInfo().name;
+          TINT_ICE(Resolver, diagnostics_) << "expected sem::VariableUser, got "
+                                           << sem_expr->TypeInfo().name;
         }
         sem_var->AddUser(sem_user);
       }
@@ -3184,7 +3131,7 @@
   }
   if (auto* vec = ty->As<sem::Vector>()) {
     if (vec->size() < 2 || vec->size() > 4) {
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Resolver, diagnostics_)
           << "Invalid vector size: vec" << vec->size();
       return false;
     }
@@ -3195,7 +3142,7 @@
   if (auto* mat = ty->As<sem::Matrix>()) {
     if (mat->columns() < 2 || mat->columns() > 4 || mat->rows() < 2 ||
         mat->rows() > 4) {
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Resolver, diagnostics_)
           << "Invalid matrix size: mat" << mat->columns() << "x" << mat->rows();
       return false;
     }
@@ -3216,7 +3163,8 @@
   if (auto* a = ty->As<sem::Atomic>()) {
     return DefaultAlignAndSize(a->Type(), align, size);
   }
-  TINT_UNREACHABLE(diagnostics_) << "invalid type " << ty->TypeInfo().name;
+  TINT_UNREACHABLE(Resolver, diagnostics_)
+      << "invalid type " << ty->TypeInfo().name;
   return false;
 }
 
@@ -3229,10 +3177,9 @@
   }
 
   if (!IsPlain(el_ty)) {  // Check must come before DefaultAlignAndSize()
-    builder_->Diagnostics().add_error(
-        el_ty->FriendlyName(builder_->Symbols()) +
-            " cannot be used as an element type of an array",
-        source);
+    AddError(el_ty->FriendlyName(builder_->Symbols()) +
+                 " cannot be used as an element type of an array",
+             source);
     return nullptr;
   }
 
@@ -3258,8 +3205,7 @@
       continue;
     }
 
-    diagnostics_.add_error("decoration is not valid for array types",
-                           deco->source());
+    AddError("decoration is not valid for array types", deco->source());
     return nullptr;
   }
 
@@ -3290,7 +3236,7 @@
       // A structure type with the block attribute must not be:
       // * the element type of an array type
       // * the member type in another structure
-      diagnostics_.add_error(
+      AddError(
           "A structure type with a [[block]] decoration cannot be used as an "
           "element of an array",
           source);
@@ -3312,7 +3258,7 @@
     // Arrays decorated with the stride attribute must have a stride that is
     // at least the size of the element type, and be a multiple of the
     // element type's alignment value.
-    diagnostics_.add_error(
+    AddError(
         "arrays decorated with the stride attribute must have a stride "
         "that is at least the size of the element type, and be a multiple "
         "of the element type's alignment value.",
@@ -3327,19 +3273,19 @@
     if (auto* r = member->Type()->As<sem::Array>()) {
       if (r->IsRuntimeSized()) {
         if (member != str->Members().back()) {
-          diagnostics_.add_error(
+          AddError(
               "v-0015",
               "runtime arrays may only appear as the last member of a struct",
               member->Declaration()->source());
           return false;
         }
         if (!str->IsBlockDecorated()) {
-          diagnostics_.add_error(
-              "v-0015",
-              "a struct containing a runtime-sized array "
-              "requires the [[block]] attribute: '" +
-                  builder_->Symbols().NameFor(str->Declaration()->name()) + "'",
-              member->Declaration()->source());
+          AddError("v-0015",
+                   "a struct containing a runtime-sized array "
+                   "requires the [[block]] attribute: '" +
+                       builder_->Symbols().NameFor(str->Declaration()->name()) +
+                       "'",
+                   member->Declaration()->source());
           return false;
         }
       }
@@ -3351,8 +3297,8 @@
             deco->Is<ast::StructMemberOffsetDecoration>() ||
             deco->Is<ast::StructMemberSizeDecoration>() ||
             deco->Is<ast::StructMemberAlignDecoration>())) {
-        diagnostics_.add_error("decoration is not valid for structure members",
-                               deco->source());
+        AddError("decoration is not valid for structure members",
+                 deco->source());
         return false;
       }
       if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
@@ -3366,11 +3312,10 @@
       if (auto* member_struct_type_block_decoration =
               ast::GetDecoration<ast::StructBlockDecoration>(
                   member_struct_type->Declaration()->decorations())) {
-        diagnostics_.add_error(
-            "structs must not contain [[block]] decorated struct members",
-            member->Declaration()->source());
-        diagnostics_.add_note("see member's struct decoration here",
-                              member_struct_type_block_decoration->source());
+        AddError("structs must not contain [[block]] decorated struct members",
+                 member->Declaration()->source());
+        AddNote("see member's struct decoration here",
+                member_struct_type_block_decoration->source());
         return false;
       }
     }
@@ -3378,8 +3323,8 @@
 
   for (auto* deco : str->Declaration()->decorations()) {
     if (!(deco->Is<ast::StructBlockDecoration>())) {
-      diagnostics_.add_error("decoration is not valid for struct declarations",
-                             deco->source());
+      AddError("decoration is not valid for struct declarations",
+               deco->source());
       return false;
     }
   }
@@ -3423,9 +3368,9 @@
 
     // Validate member type
     if (!IsPlain(type)) {
-      diagnostics_.add_error(
-          type->FriendlyName(builder_->Symbols()) +
-          " cannot be used as the type of a structure member");
+      AddError(type->FriendlyName(builder_->Symbols()) +
+                   " cannot be used as the type of a structure member",
+               member->source());
       return nullptr;
     }
 
@@ -3449,8 +3394,7 @@
         // Offset decorations are not part of the WGSL spec, but are emitted
         // by the SPIR-V reader.
         if (o->offset() < struct_size) {
-          diagnostics_.add_error("offsets must be in ascending order",
-                                 o->source());
+          AddError("offsets must be in ascending order", o->source());
           return nullptr;
         }
         offset = o->offset();
@@ -3458,19 +3402,17 @@
         has_offset_deco = true;
       } else if (auto* a = deco->As<ast::StructMemberAlignDecoration>()) {
         if (a->align() <= 0 || !utils::IsPowerOfTwo(a->align())) {
-          diagnostics_.add_error(
-              "align value must be a positive, power-of-two integer",
-              a->source());
+          AddError("align value must be a positive, power-of-two integer",
+                   a->source());
           return nullptr;
         }
         align = a->align();
         has_align_deco = true;
       } else if (auto* s = deco->As<ast::StructMemberSizeDecoration>()) {
         if (s->size() < size) {
-          diagnostics_.add_error(
-              "size must be at least as big as the type's size (" +
-                  std::to_string(size) + ")",
-              s->source());
+          AddError("size must be at least as big as the type's size (" +
+                       std::to_string(size) + ")",
+                   s->source());
           return nullptr;
         }
         size = s->size();
@@ -3479,7 +3421,7 @@
     }
 
     if (has_offset_deco && (has_align_deco || has_size_deco)) {
-      diagnostics_.add_error(
+      AddError(
           "offset decorations cannot be used with align or size decorations",
           member->source());
       return nullptr;
@@ -3525,13 +3467,12 @@
                                     : builder_->create<sem::Void>();
 
   if (func_type->UnwrapRef() != ret_type) {
-    diagnostics_.add_error("v-000y",
-                           "return statement type must match its function "
-                           "return type, returned '" +
-                               ret_type->FriendlyName(builder_->Symbols()) +
-                               "', expected '" +
-                               current_function_->return_type_name + "'",
-                           ret->source());
+    AddError("v-000y",
+             "return statement type must match its function "
+             "return type, returned '" +
+                 ret_type->FriendlyName(builder_->Symbols()) + "', expected '" +
+                 current_function_->return_type_name + "'",
+             ret->source());
     return false;
   }
 
@@ -3556,10 +3497,10 @@
 bool Resolver::ValidateSwitch(const ast::SwitchStatement* s) {
   auto* cond_type = TypeOf(s->condition())->UnwrapRef();
   if (!cond_type->is_integer_scalar()) {
-    diagnostics_.add_error("v-0025",
-                           "switch statement selector expression must be of a "
-                           "scalar integer type",
-                           s->condition()->source());
+    AddError("v-0025",
+             "switch statement selector expression must be of a "
+             "scalar integer type",
+             s->condition()->source());
     return false;
   }
 
@@ -3570,9 +3511,9 @@
     if (case_stmt->IsDefault()) {
       if (has_default) {
         // More than one default clause
-        diagnostics_.add_error(
-            "v-0008", "switch statement must have exactly one default clause",
-            case_stmt->source());
+        AddError("v-0008",
+                 "switch statement must have exactly one default clause",
+                 case_stmt->source());
         return false;
       }
       has_default = true;
@@ -3580,21 +3521,20 @@
 
     for (auto* selector : case_stmt->selectors()) {
       if (cond_type != TypeOf(selector)) {
-        diagnostics_.add_error("v-0026",
-                               "the case selector values must have the same "
-                               "type as the selector expression.",
-                               case_stmt->source());
+        AddError("v-0026",
+                 "the case selector values must have the same "
+                 "type as the selector expression.",
+                 case_stmt->source());
         return false;
       }
 
       auto v = selector->value_as_u32();
       if (selector_set.find(v) != selector_set.end()) {
-        diagnostics_.add_error(
-            "v-0027",
-            "a literal value must not appear more than once in "
-            "the case selectors for a switch statement: '" +
-                builder_->str(selector) + "'",
-            case_stmt->source());
+        AddError("v-0027",
+                 "a literal value must not appear more than once in "
+                 "the case selectors for a switch statement: '" +
+                     builder_->str(selector) + "'",
+                 case_stmt->source());
         return false;
       }
       selector_set.emplace(v);
@@ -3603,8 +3543,7 @@
 
   if (!has_default) {
     // No default clause
-    diagnostics_.add_error("switch statement must have a default clause",
-                           s->source());
+    AddError("switch statement must have a default clause", s->source());
     return false;
   }
 
@@ -3612,10 +3551,10 @@
     auto* last_clause = s->body().back()->As<ast::CaseStatement>();
     auto* last_stmt = last_clause->body()->last();
     if (last_stmt && last_stmt->Is<ast::FallthroughStatement>()) {
-      diagnostics_.add_error("v-0028",
-                             "a fallthrough statement must not appear as "
-                             "the last statement in last clause of a switch",
-                             last_stmt->source());
+      AddError("v-0028",
+               "a fallthrough statement must not appear as "
+               "the last statement in last clause of a switch",
+               last_stmt->source());
       return false;
     }
   }
@@ -3663,9 +3602,8 @@
   auto* lhs_ref = lhs_type->As<sem::Reference>();
   if (!lhs_ref) {
     // LHS is not a reference, so it has no storage.
-    diagnostics_.add_error(
-        "cannot assign to value of type '" + TypeNameOf(a->lhs()) + "'",
-        a->lhs()->source());
+    AddError("cannot assign to value of type '" + TypeNameOf(a->lhs()) + "'",
+             a->lhs()->source());
 
     return false;
   }
@@ -3675,13 +3613,13 @@
 
   // Value type has to match storage type
   if (storage_type != value_type) {
-    diagnostics_.add_error("cannot assign '" + TypeNameOf(a->rhs()) + "' to '" +
-                               TypeNameOf(a->lhs()) + "'",
-                           a->source());
+    AddError("cannot assign '" + TypeNameOf(a->rhs()) + "' to '" +
+                 TypeNameOf(a->lhs()) + "'",
+             a->source());
     return false;
   }
   if (lhs_ref->Access() == ast::Access::kRead) {
-    diagnostics_.add_error(
+    AddError(
         "cannot store into a read-only type '" + TypeNameOf(a->lhs()) + "'",
         a->source());
     return false;
@@ -3695,10 +3633,8 @@
   for (auto* d : decorations) {
     auto res = seen.emplace(&d->TypeInfo(), d->source());
     if (!res.second && !d->Is<ast::InternalDecoration>()) {
-      diagnostics_.add_error("duplicate " + d->name() + " decoration",
-                             d->source());
-      diagnostics_.add_note("first decoration declared here",
-                            res.first->second);
+      AddError("duplicate " + d->name() + " decoration", d->source());
+      AddNote("first decoration declared here", res.first->second);
       return false;
     }
   }
@@ -3723,7 +3659,7 @@
         err << "while analysing structure member "
             << str->FriendlyName(builder_->Symbols()) << "."
             << builder_->Symbols().NameFor(member->Declaration()->symbol());
-        diagnostics_.add_note(err.str(), member->Declaration()->source());
+        AddNote(err.str(), member->Declaration()->source());
         return false;
       }
     }
@@ -3740,7 +3676,7 @@
     err << "Type '" << ty->FriendlyName(builder_->Symbols())
         << "' cannot be used in storage class '" << sc
         << "' as it is non-host-shareable";
-    diagnostics_.add_error(err.str(), usage);
+    AddError(err.str(), usage);
     return false;
   }
 
@@ -3758,7 +3694,7 @@
       // Recurse into the constructor argument expression.
       return GetScalarConstExprValue(type_constructor->values()[0], result);
     } else {
-      TINT_ICE(diagnostics_) << "malformed scalar type constructor";
+      TINT_ICE(Resolver, diagnostics_) << "malformed scalar type constructor";
     }
   } else if (auto* scalar = expr->As<ast::ScalarConstructorExpression>()) {
     // Cast literal to result type.
@@ -3772,10 +3708,10 @@
       *result = static_cast<T>(bool_lit->IsTrue());
       return true;
     } else {
-      TINT_ICE(diagnostics_) << "unhandled scalar constructor";
+      TINT_ICE(Resolver, diagnostics_) << "unhandled scalar constructor";
     }
   } else {
-    TINT_ICE(diagnostics_) << "unhandled constant expression";
+    TINT_ICE(Resolver, diagnostics_) << "unhandled constant expression";
   }
 
   return false;
@@ -3785,8 +3721,9 @@
 bool Resolver::BlockScope(const ast::BlockStatement* block, F&& callback) {
   auto* sem_block = builder_->Sem().Get<sem::BlockStatement>(block);
   if (!sem_block) {
-    TINT_ICE(diagnostics_) << "Resolver::BlockScope() called on a block for "
-                              "which semantic information is not available";
+    TINT_ICE(Resolver, diagnostics_)
+        << "Resolver::BlockScope() called on a block for "
+           "which semantic information is not available";
     return false;
   }
   TINT_SCOPED_ASSIGNMENT(current_block_,
@@ -3805,12 +3742,12 @@
 
 void Resolver::Mark(const ast::Node* node) {
   if (node == nullptr) {
-    TINT_ICE(diagnostics_) << "Resolver::Mark() called with nullptr";
+    TINT_ICE(Resolver, diagnostics_) << "Resolver::Mark() called with nullptr";
   }
   if (marked_.emplace(node).second) {
     return;
   }
-  TINT_ICE(diagnostics_)
+  TINT_ICE(Resolver, diagnostics_)
       << "AST node '" << node->TypeInfo().name
       << "' was encountered twice in the same AST of a Program\n"
       << "At: " << node->source() << "\n"
@@ -3818,6 +3755,24 @@
       << "Pointer: " << node;
 }
 
+void Resolver::AddError(const char* code,
+                        const std::string& msg,
+                        const Source& source) const {
+  diagnostics_.add_error(diag::System::Resolver, code, msg, source);
+}
+
+void Resolver::AddError(const std::string& msg, const Source& source) const {
+  diagnostics_.add_error(diag::System::Resolver, msg, source);
+}
+
+void Resolver::AddWarning(const std::string& msg, const Source& source) const {
+  diagnostics_.add_warning(diag::System::Resolver, msg, source);
+}
+
+void Resolver::AddNote(const std::string& msg, const Source& source) const {
+  diagnostics_.add_note(diag::System::Resolver, msg, source);
+}
+
 Resolver::VariableInfo::VariableInfo(const ast::Variable* decl,
                                      sem::Type* ty,
                                      const std::string& tn,
diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h
index cf303c0..531912f 100644
--- a/src/resolver/resolver.h
+++ b/src/resolver/resolver.h
@@ -409,6 +409,21 @@
   /// @param node the AST node.
   void Mark(const ast::Node* node);
 
+  /// Adds the given error message to the diagnostics
+  /// [DEPRECATED] Remove all codes
+  void AddError(const char* code,
+                const std::string& msg,
+                const Source& source) const;
+
+  /// Adds the given error message to the diagnostics
+  void AddError(const std::string& msg, const Source& source) const;
+
+  /// Adds the given warning message to the diagnostics
+  void AddWarning(const std::string& msg, const Source& source) const;
+
+  /// Adds the given note message to the diagnostics
+  void AddNote(const std::string& msg, const Source& source) const;
+
   template <typename CALLBACK>
   void TraverseCallChain(FunctionInfo* from,
                          FunctionInfo* to,
diff --git a/src/sem/array.cc b/src/sem/array.cc
index 5520bb5..9dd6ea8 100644
--- a/src/sem/array.cc
+++ b/src/sem/array.cc
@@ -35,7 +35,7 @@
       size_(size),
       stride_(stride),
       implicit_stride_(implicit_stride) {
-  TINT_ASSERT(element_);
+  TINT_ASSERT(Semantic, element_);
 }
 
 std::string Array::type_name() const {
diff --git a/src/sem/atomic_type.cc b/src/sem/atomic_type.cc
index e632bd8..4444768 100644
--- a/src/sem/atomic_type.cc
+++ b/src/sem/atomic_type.cc
@@ -23,7 +23,7 @@
 namespace sem {
 
 Atomic::Atomic(const sem::Type* subtype) : subtype_(subtype) {
-  TINT_ASSERT(!subtype->Is<Reference>());
+  TINT_ASSERT(AST, !subtype->Is<Reference>());
 }
 
 std::string Atomic::type_name() const {
diff --git a/src/sem/call_target.cc b/src/sem/call_target.cc
index d8c5f49..ce4c828 100644
--- a/src/sem/call_target.cc
+++ b/src/sem/call_target.cc
@@ -23,7 +23,7 @@
 
 CallTarget::CallTarget(sem::Type* return_type, const ParameterList& parameters)
     : return_type_(return_type), parameters_(parameters) {
-  TINT_ASSERT(return_type);
+  TINT_ASSERT(Semantic, return_type);
 }
 
 CallTarget::~CallTarget() = default;
diff --git a/src/sem/depth_texture_type.cc b/src/sem/depth_texture_type.cc
index d028d68..f9cd3ec 100644
--- a/src/sem/depth_texture_type.cc
+++ b/src/sem/depth_texture_type.cc
@@ -32,7 +32,7 @@
 }  // namespace
 
 DepthTexture::DepthTexture(ast::TextureDimension dim) : Base(dim) {
-  TINT_ASSERT(IsValidDepthDimension(dim));
+  TINT_ASSERT(Semantic, IsValidDepthDimension(dim));
 }
 
 DepthTexture::DepthTexture(DepthTexture&&) = default;
diff --git a/src/sem/expression.cc b/src/sem/expression.cc
index 7286dc6..b5ee7f4 100644
--- a/src/sem/expression.cc
+++ b/src/sem/expression.cc
@@ -23,7 +23,7 @@
                        const sem::Type* type,
                        Statement* statement)
     : declaration_(declaration), type_(type), statement_(statement) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 }  // namespace sem
diff --git a/src/sem/info.h b/src/sem/info.h
index dfb26f2..78b82d5 100644
--- a/src/sem/info.h
+++ b/src/sem/info.h
@@ -70,7 +70,7 @@
   void Add(const AST_OR_TYPE* node,
            const SemanticNodeTypeFor<AST_OR_TYPE>* sem_node) {
     // Check there's no semantic info already existing for the node
-    TINT_ASSERT(Get(node) == nullptr);
+    TINT_ASSERT(Semantic, Get(node) == nullptr);
     map.emplace(node, sem_node);
   }
 
diff --git a/src/sem/matrix_type.cc b/src/sem/matrix_type.cc
index 06e392e..2de6a01 100644
--- a/src/sem/matrix_type.cc
+++ b/src/sem/matrix_type.cc
@@ -27,10 +27,10 @@
       column_type_(column_type),
       rows_(column_type->size()),
       columns_(columns) {
-  TINT_ASSERT(rows_ > 1);
-  TINT_ASSERT(rows_ < 5);
-  TINT_ASSERT(columns_ > 1);
-  TINT_ASSERT(columns_ < 5);
+  TINT_ASSERT(AST, rows_ > 1);
+  TINT_ASSERT(AST, rows_ < 5);
+  TINT_ASSERT(AST, columns_ > 1);
+  TINT_ASSERT(AST, columns_ < 5);
 }
 
 Matrix::Matrix(Matrix&&) = default;
diff --git a/src/sem/multisampled_texture_type.cc b/src/sem/multisampled_texture_type.cc
index ad84dc7..ad9cc67 100644
--- a/src/sem/multisampled_texture_type.cc
+++ b/src/sem/multisampled_texture_type.cc
@@ -24,7 +24,7 @@
 MultisampledTexture::MultisampledTexture(ast::TextureDimension dim,
                                          const Type* type)
     : Base(dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 MultisampledTexture::MultisampledTexture(MultisampledTexture&&) = default;
diff --git a/src/sem/pointer_type.cc b/src/sem/pointer_type.cc
index d4b9734..1ab5d1a 100644
--- a/src/sem/pointer_type.cc
+++ b/src/sem/pointer_type.cc
@@ -26,8 +26,8 @@
                  ast::StorageClass storage_class,
                  ast::Access access)
     : subtype_(subtype), storage_class_(storage_class), access_(access) {
-  TINT_ASSERT(!subtype->Is<Reference>());
-  TINT_ASSERT(access != ast::Access::kUndefined);
+  TINT_ASSERT(Semantic, !subtype->Is<Reference>());
+  TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
 std::string Pointer::type_name() const {
diff --git a/src/sem/reference_type.cc b/src/sem/reference_type.cc
index 959c61d..48bde77 100644
--- a/src/sem/reference_type.cc
+++ b/src/sem/reference_type.cc
@@ -25,8 +25,8 @@
                      ast::StorageClass storage_class,
                      ast::Access access)
     : subtype_(subtype), storage_class_(storage_class), access_(access) {
-  TINT_ASSERT(!subtype->Is<Reference>());
-  TINT_ASSERT(access != ast::Access::kUndefined);
+  TINT_ASSERT(Semantic, !subtype->Is<Reference>());
+  TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
 std::string Reference::type_name() const {
diff --git a/src/sem/sampled_texture_type.cc b/src/sem/sampled_texture_type.cc
index c54ef1b..21a30a4 100644
--- a/src/sem/sampled_texture_type.cc
+++ b/src/sem/sampled_texture_type.cc
@@ -23,7 +23,7 @@
 
 SampledTexture::SampledTexture(ast::TextureDimension dim, const Type* type)
     : Base(dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 SampledTexture::SampledTexture(SampledTexture&&) = default;
diff --git a/src/sem/statement.cc b/src/sem/statement.cc
index 3988710..c7d46fb 100644
--- a/src/sem/statement.cc
+++ b/src/sem/statement.cc
@@ -47,7 +47,7 @@
             }
           }
         }
-        TINT_ASSERT(statement_is_continuing_for_loop);
+        TINT_ASSERT(Semantic, statement_is_continuing_for_loop);
       }
     }
   }
diff --git a/src/sem/vector_type.cc b/src/sem/vector_type.cc
index 9b4ef83..4b6fe20 100644
--- a/src/sem/vector_type.cc
+++ b/src/sem/vector_type.cc
@@ -23,8 +23,8 @@
 
 Vector::Vector(Type const* subtype, uint32_t size)
     : subtype_(subtype), size_(size) {
-  TINT_ASSERT(size_ > 1);
-  TINT_ASSERT(size_ < 5);
+  TINT_ASSERT(Semantic, size_ > 1);
+  TINT_ASSERT(Semantic, size_ < 5);
 }
 
 Vector::Vector(Vector&&) = default;
diff --git a/src/symbol.cc b/src/symbol.cc
index 13db168..e77c699 100644
--- a/src/symbol.cc
+++ b/src/symbol.cc
@@ -32,7 +32,8 @@
 Symbol& Symbol::operator=(Symbol&& o) = default;
 
 bool Symbol::operator==(const Symbol& other) const {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(program_id_, other.program_id_);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Symbol, program_id_,
+                                         other.program_id_);
   return val_ == other.val_;
 }
 
diff --git a/src/symbol_table.cc b/src/symbol_table.cc
index a667e7d..f5a2655 100644
--- a/src/symbol_table.cc
+++ b/src/symbol_table.cc
@@ -54,7 +54,7 @@
 }
 
 std::string SymbolTable::NameFor(const Symbol symbol) const {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL(program_id_, symbol);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL(Symbol, program_id_, symbol);
   auto it = symbol_to_name_.find(symbol);
   if (it == symbol_to_name_.end()) {
     return symbol.to_str();
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index e965cae..e7fa0f7 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -38,6 +38,7 @@
   auto* cfg = data.Get<Config>();
   if (cfg == nullptr) {
     out.Diagnostics().add_error(
+        diag::System::Transform,
         "missing transform data for ArrayLengthFromUniform");
     return Output(Program(std::move(out)));
   }
@@ -95,13 +96,13 @@
     // have been run before this transform.
     auto* param = call_expr->params()[0]->As<ast::UnaryOpExpression>();
     if (!param || param->op() != ast::UnaryOp::kAddressOf) {
-      TINT_ICE(ctx.dst->Diagnostics())
+      TINT_ICE(Transform, ctx.dst->Diagnostics())
           << "expected form of arrayLength argument to be &resource.array";
       break;
     }
     auto* accessor = param->expr()->As<ast::MemberAccessorExpression>();
     if (!accessor) {
-      TINT_ICE(ctx.dst->Diagnostics())
+      TINT_ICE(Transform, ctx.dst->Diagnostics())
           << "expected form of arrayLength argument to be &resource.array";
       break;
     }
@@ -109,7 +110,7 @@
     auto* storage_buffer_sem =
         sem.Get(storage_buffer_expr)->As<sem::VariableUser>();
     if (!storage_buffer_sem) {
-      TINT_ICE(ctx.dst->Diagnostics())
+      TINT_ICE(Transform, ctx.dst->Diagnostics())
           << "expected form of arrayLength argument to be &resource.array";
       break;
     }
@@ -119,9 +120,10 @@
     auto idx_itr = cfg->bindpoint_to_size_index.find(binding);
     if (idx_itr == cfg->bindpoint_to_size_index.end()) {
       ctx.dst->Diagnostics().add_error(
+          diag::System::Transform,
           "missing size index mapping for binding point (" +
-          std::to_string(binding.group) + "," +
-          std::to_string(binding.binding) + ")");
+              std::to_string(binding.group) + "," +
+              std::to_string(binding.binding) + ")");
       continue;
     }
 
diff --git a/src/transform/binding_remapper.cc b/src/transform/binding_remapper.cc
index 3755fff..5873cfe 100644
--- a/src/transform/binding_remapper.cc
+++ b/src/transform/binding_remapper.cc
@@ -45,6 +45,7 @@
   auto* remappings = datamap.Get<Remappings>();
   if (!remappings) {
     out.Diagnostics().add_error(
+        diag::System::Transform,
         "BindingRemapper did not find the remapping data");
     return Output(Program(std::move(out)));
   }
diff --git a/src/transform/bound_array_accessors.cc b/src/transform/bound_array_accessors.cc
index 8d7b68e..e2b1860 100644
--- a/src/transform/bound_array_accessors.cc
+++ b/src/transform/bound_array_accessors.cc
@@ -74,7 +74,8 @@
       auto* limit = b.Sub(arr_len, b.Expr(1u));
       new_idx = b.Call("min", b.Construct<u32>(ctx->Clone(old_idx)), limit);
     } else {
-      diags.add_error("invalid 0 size", expr->source());
+      diags.add_error(diag::System::Transform, "invalid 0 size",
+                      expr->source());
       return nullptr;
     }
   } else if (auto* c = old_idx->As<ast::ScalarConstructorExpression>()) {
@@ -86,7 +87,8 @@
     } else if (auto* uint = lit->As<ast::UintLiteral>()) {
       new_idx = b.Expr(std::min(uint->value(), size - 1));
     } else {
-      diags.add_error("unknown scalar constructor type for accessor",
+      diags.add_error(diag::System::Transform,
+                      "unknown scalar constructor type for accessor",
                       expr->source());
       return nullptr;
     }
diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc
index 4a6f571..ee48611 100644
--- a/src/transform/calculate_array_length.cc
+++ b/src/transform/calculate_array_length.cc
@@ -134,7 +134,7 @@
           auto* arg = call_expr->params()[0];
           auto* address_of = arg->As<ast::UnaryOpExpression>();
           if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) {
-            TINT_ICE(ctx.dst->Diagnostics())
+            TINT_ICE(Transform, ctx.dst->Diagnostics())
                 << "arrayLength() expected pointer to member access, got "
                 << address_of->TypeInfo().name;
           }
@@ -142,7 +142,7 @@
 
           auto* accessor = array_expr->As<ast::MemberAccessorExpression>();
           if (!accessor) {
-            TINT_ICE(ctx.dst->Diagnostics())
+            TINT_ICE(Transform, ctx.dst->Diagnostics())
                 << "arrayLength() expected pointer to member access, got "
                    "pointer to "
                 << array_expr->TypeInfo().name;
@@ -158,7 +158,7 @@
           auto buffer_size = get_buffer_size_intrinsic(storage_buffer_type);
 
           if (!storage_buffer_type) {
-            TINT_ICE(ctx.dst->Diagnostics())
+            TINT_ICE(Transform, ctx.dst->Diagnostics())
                 << "arrayLength(X.Y) expected X to be sem::Struct, got "
                 << storage_buffer_type->FriendlyName(ctx.src->Symbols());
             break;
diff --git a/src/transform/canonicalize_entry_point_io.cc b/src/transform/canonicalize_entry_point_io.cc
index 41fc8fa..ebac0b9 100644
--- a/src/transform/canonicalize_entry_point_io.cc
+++ b/src/transform/canonicalize_entry_point_io.cc
@@ -69,6 +69,7 @@
   auto* cfg = data.Get<Config>();
   if (cfg == nullptr) {
     out.Diagnostics().add_error(
+        diag::System::Transform,
         "missing transform data for CanonicalizeEntryPointIO");
     return Output(Program(std::move(out)));
   }
@@ -146,7 +147,8 @@
           ast::ExpressionList init_values;
           for (auto* member : str->Members()) {
             if (member->Type()->Is<sem::Struct>()) {
-              TINT_ICE(ctx.dst->Diagnostics()) << "nested pipeline IO struct";
+              TINT_ICE(Transform, ctx.dst->Diagnostics())
+                  << "nested pipeline IO struct";
             }
 
             if (cfg->builtin_style == BuiltinStyle::kParameter &&
@@ -239,7 +241,8 @@
         // Rebuild struct with only the entry point IO attributes.
         for (auto* member : str->Members()) {
           if (member->Type()->Is<sem::Struct>()) {
-            TINT_ICE(ctx.dst->Diagnostics()) << "nested pipeline IO struct";
+            TINT_ICE(Transform, ctx.dst->Diagnostics())
+                << "nested pipeline IO struct";
           }
 
           ast::DecorationList new_decorations = RemoveDecorations(
diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index 57e40346..96a373a 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -352,7 +352,7 @@
       op = DecomposeMemoryAccess::Intrinsic::Op::kAtomicCompareExchangeWeak;
       break;
     default:
-      TINT_ICE(builder->Diagnostics())
+      TINT_ICE(Transform, builder->Diagnostics())
           << "invalid IntrinsicType for DecomposeMemoryAccess::Intrinsic: "
           << ty->type_name();
       break;
@@ -435,7 +435,7 @@
   /// @param expr the expression that performs the access
   /// @param access the access
   void AddAccess(ast::Expression* expr, BufferAccess&& access) {
-    TINT_ASSERT(access.type);
+    TINT_ASSERT(Transform, access.type);
     accesses.emplace(expr, std::move(access));
     expression_order.emplace_back(expr);
   }
@@ -672,7 +672,7 @@
 
       auto* atomic = IntrinsicAtomicFor(ctx.dst, op, el_ty);
       if (atomic == nullptr) {
-        TINT_ICE(ctx.dst->Diagnostics())
+        TINT_ICE(Transform, ctx.dst->Diagnostics())
             << "IntrinsicAtomicFor() returned nullptr for op " << op
             << " and type " << el_ty->type_name();
       }
diff --git a/src/transform/external_texture_transform.cc b/src/transform/external_texture_transform.cc
index 81e4fab..4d48ff7 100644
--- a/src/transform/external_texture_transform.cc
+++ b/src/transform/external_texture_transform.cc
@@ -58,7 +58,7 @@
                     ->Is<sem::ExternalTexture>()) {
               if (intrinsic->Type() == sem::IntrinsicType::kTextureLoad &&
                   call_expr->params().size() != 2) {
-                TINT_ICE(ctx.dst->Diagnostics())
+                TINT_ICE(Transform, ctx.dst->Diagnostics())
                     << "expected textureLoad call with a texture_external to "
                        "have 2 parameters, found "
                     << call_expr->params().size() << " parameters";
@@ -67,7 +67,7 @@
               if (intrinsic->Type() ==
                       sem::IntrinsicType::kTextureSampleLevel &&
                   call_expr->params().size() != 3) {
-                TINT_ICE(ctx.dst->Diagnostics())
+                TINT_ICE(Transform, ctx.dst->Diagnostics())
                     << "expected textureSampleLevel call with a "
                        "texture_external to have 3 parameters, found "
                     << call_expr->params().size() << " parameters";
diff --git a/src/transform/fold_constants.cc b/src/transform/fold_constants.cc
index 9ef726b..3fe3691 100644
--- a/src/transform/fold_constants.cc
+++ b/src/transform/fold_constants.cc
@@ -67,7 +67,7 @@
   operator bool() const { return Valid(); }
 
   void Append(const Value& value) {
-    TINT_ASSERT(value.type == type);
+    TINT_ASSERT(Transform, value.type == type);
     elems.insert(elems.end(), value.elems.begin(), value.elems.end());
   }
 
@@ -89,7 +89,7 @@
         return func(elems[index].bool_);
       }
     }
-    TINT_ASSERT(false && "Unreachable");
+    TINT_ASSERT(Transform, false && "Unreachable");
     return func(~0);
   }
 };
@@ -105,7 +105,7 @@
   } else if (t->Is<ast::Bool>()) {
     return Value::Type::bool_;
   }
-  TINT_ASSERT(false && "Invalid type");
+  TINT_ASSERT(Transform, false && "Invalid type");
   return {};
 }
 
@@ -193,7 +193,7 @@
   if (auto* lit = literal->As<ast::BoolLiteral>()) {
     return {lit->IsTrue()};
   }
-  TINT_ASSERT(false && "Unreachable");
+  TINT_ASSERT(Transform, false && "Unreachable");
   return {};
 }
 
diff --git a/src/transform/promote_initializers_to_const_var.cc b/src/transform/promote_initializers_to_const_var.cc
index aed8e5d..d0724d3 100644
--- a/src/transform/promote_initializers_to_const_var.cc
+++ b/src/transform/promote_initializers_to_const_var.cc
@@ -52,7 +52,7 @@
     if (auto* src_init = src_node->As<ast::TypeConstructorExpression>()) {
       auto* src_sem_expr = ctx.src->Sem().Get(src_init);
       if (!src_sem_expr) {
-        TINT_ICE(ctx.dst->Diagnostics())
+        TINT_ICE(Transform, ctx.dst->Diagnostics())
             << "ast::TypeConstructorExpression has no semantic expression node";
         continue;
       }
diff --git a/src/transform/renamer.cc b/src/transform/renamer.cc
index d1371b2..c39171f 100644
--- a/src/transform/renamer.cc
+++ b/src/transform/renamer.cc
@@ -857,7 +857,7 @@
     if (auto* member = node->As<ast::MemberAccessorExpression>()) {
       auto* sem = in->Sem().Get(member);
       if (!sem) {
-        TINT_ICE(out.Diagnostics())
+        TINT_ICE(Transform, out.Diagnostics())
             << "MemberAccessorExpression has no semantic info";
         continue;
       }
@@ -867,7 +867,8 @@
     } else if (auto* call = node->As<ast::CallExpression>()) {
       auto* sem = in->Sem().Get(call);
       if (!sem) {
-        TINT_ICE(out.Diagnostics()) << "CallExpression has no semantic info";
+        TINT_ICE(Transform, out.Diagnostics())
+            << "CallExpression has no semantic info";
         continue;
       }
       if (sem->Target()->Is<sem::Intrinsic>()) {
diff --git a/src/transform/single_entry_point.cc b/src/transform/single_entry_point.cc
index a926d11..15afa5f 100644
--- a/src/transform/single_entry_point.cc
+++ b/src/transform/single_entry_point.cc
@@ -35,7 +35,8 @@
 
   auto* cfg = data.Get<Config>();
   if (cfg == nullptr) {
-    out.Diagnostics().add_error("missing transform data for SingleEntryPoint");
+    out.Diagnostics().add_error(diag::System::Transform,
+                                "missing transform data for SingleEntryPoint");
     return Output(Program(std::move(out)));
   }
 
@@ -51,8 +52,9 @@
     }
   }
   if (entry_point == nullptr) {
-    out.Diagnostics().add_error("entry point '" + cfg->entry_point_name +
-                                "' not found");
+    out.Diagnostics().add_error(
+        diag::System::Transform,
+        "entry point '" + cfg->entry_point_name + "' not found");
     return Output(Program(std::move(out)));
   }
 
@@ -81,7 +83,7 @@
         out.AST().AddFunction(ctx.Clone(func));
       }
     } else {
-      TINT_UNREACHABLE(out.Diagnostics())
+      TINT_UNREACHABLE(Transform, out.Diagnostics())
           << "unhandled global declaration: " << decl->TypeInfo().name;
       return Output(Program(std::move(out)));
     }
diff --git a/src/transform/transform.cc b/src/transform/transform.cc
index c334337..914b167 100644
--- a/src/transform/transform.cc
+++ b/src/transform/transform.cc
@@ -134,7 +134,7 @@
   if (auto* s = ty->As<sem::Sampler>()) {
     return ctx->dst->create<ast::Sampler>(s->kind());
   }
-  TINT_UNREACHABLE(ctx->dst->Diagnostics())
+  TINT_UNREACHABLE(Transform, ctx->dst->Diagnostics())
       << "Unhandled type: " << ty->TypeInfo().name;
   return nullptr;
 }
diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc
index dc7e70c..e06dd7c 100644
--- a/src/transform/vertex_pulling.cc
+++ b/src/transform/vertex_pulling.cc
@@ -304,7 +304,8 @@
       }
       new_function_parameters.push_back(ctx.Clone(param));
     } else {
-      TINT_ICE(ctx.dst->Diagnostics()) << "Invalid entry point parameter";
+      TINT_ICE(Transform, ctx.dst->Diagnostics())
+          << "Invalid entry point parameter";
     }
   }
 
@@ -346,7 +347,8 @@
         }
         members_to_clone.push_back(member);
       } else {
-        TINT_ICE(ctx.dst->Diagnostics()) << "Invalid entry point parameter";
+        TINT_ICE(Transform, ctx.dst->Diagnostics())
+            << "Invalid entry point parameter";
       }
     }
 
@@ -466,7 +468,8 @@
   auto* func = in->AST().Functions().Find(
       in->Symbols().Get(cfg.entry_point_name), ast::PipelineStage::kVertex);
   if (func == nullptr) {
-    out.Diagnostics().add_error("Vertex stage entry point not found");
+    out.Diagnostics().add_error(diag::System::Transform,
+                                "Vertex stage entry point not found");
     return Output(Program(std::move(out)));
   }
 
diff --git a/src/transform/zero_init_workgroup_memory.cc b/src/transform/zero_init_workgroup_memory.cc
index 100a379..287018c 100644
--- a/src/transform/zero_init_workgroup_memory.cc
+++ b/src/transform/zero_init_workgroup_memory.cc
@@ -78,7 +78,7 @@
       return;
     }
 
-    TINT_UNREACHABLE(ctx.dst->Diagnostics())
+    TINT_UNREACHABLE(Transform, ctx.dst->Diagnostics())
         << "could not zero workgroup type: " << ty->type_name();
   }
 
diff --git a/src/utils/io/tmpfile_posix.cc b/src/utils/io/tmpfile_posix.cc
index 9d66331..39d2af9 100644
--- a/src/utils/io/tmpfile_posix.cc
+++ b/src/utils/io/tmpfile_posix.cc
@@ -30,8 +30,8 @@
   // (when the source value exceeds the representable range) is implementation
   // defined. While such a large file extension is unlikely in practice, we
   // enforce this here at runtime.
-  TINT_ASSERT(ext.length() <=
-              static_cast<size_t>(std::numeric_limits<int>::max()));
+  TINT_ASSERT(Utils, ext.length() <=
+                         static_cast<size_t>(std::numeric_limits<int>::max()));
   std::string name = "tint_XXXXXX" + ext;
   int file = mkstemps(&name[0], static_cast<int>(ext.length()));
   if (file != -1) {
diff --git a/src/writer/append_vector.cc b/src/writer/append_vector.cc
index e61ffe4..9763c18 100644
--- a/src/writer/append_vector.cc
+++ b/src/writer/append_vector.cc
@@ -60,8 +60,9 @@
   } else if (packed_el_sem_ty->Is<sem::Bool>()) {
     packed_el_ty = b->create<ast::Bool>();
   } else {
-    TINT_UNREACHABLE(b->Diagnostics()) << "unsupported vector element type: "
-                                       << packed_el_sem_ty->TypeInfo().name;
+    TINT_UNREACHABLE(Writer, b->Diagnostics())
+        << "unsupported vector element type: "
+        << packed_el_sem_ty->TypeInfo().name;
   }
 
   auto* statement = vector_sem->Stmt();
@@ -86,7 +87,7 @@
         } else if (packed_el_sem_ty->Is<sem::Bool>()) {
           return b->Expr(false);
         } else {
-          TINT_UNREACHABLE(b->Diagnostics())
+          TINT_UNREACHABLE(Writer, b->Diagnostics())
               << "unsupported vector element type: "
               << packed_el_sem_ty->TypeInfo().name;
         }
diff --git a/src/writer/float_to_string.cc b/src/writer/float_to_string.cc
index 1692f26..81e66bc 100644
--- a/src/writer/float_to_string.cc
+++ b/src/writer/float_to_string.cc
@@ -106,7 +106,7 @@
         }
       } else {
         // Subnormal, and not zero.
-        TINT_ASSERT(mantissa != 0);
+        TINT_ASSERT(Writer, mantissa != 0);
         const int kTopBit = (1 << kMantissaBits);
 
         // Shift left until we get 1.x
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index b2b4c502..ce1aea8 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -152,7 +152,7 @@
         }
       }
     } else {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Writer, diagnostics_)
           << "unhandled module-scope declaration: " << decl->TypeInfo().name;
       return false;
     }
@@ -188,7 +188,8 @@
                                 ast::BitcastExpression* expr) {
   auto* type = TypeOf(expr);
   if (!type->is_integer_scalar() && !type->is_float_scalar()) {
-    diagnostics_.add_error("Unable to do bitcast to type " + type->type_name());
+    diagnostics_.add_error(diag::System::Writer,
+                           "Unable to do bitcast to type " + type->type_name());
     return false;
   }
 
@@ -305,7 +306,7 @@
     case ast::BinaryOp::kLogicalAnd:
     case ast::BinaryOp::kLogicalOr: {
       // These are both handled above.
-      TINT_UNREACHABLE(diagnostics_);
+      TINT_UNREACHABLE(Writer, diagnostics_);
       return false;
     }
     case ast::BinaryOp::kEqual:
@@ -353,7 +354,8 @@
       out << "%";
       break;
     case ast::BinaryOp::kNone:
-      diagnostics_.add_error("missing binary operation type");
+      diagnostics_.add_error(diag::System::Writer,
+                             "missing binary operation type");
       return false;
   }
   out << " ";
@@ -437,7 +439,7 @@
         case ast::StorageClass::kStorage:
           return EmitStorageBufferAccess(pre, out, expr, intrinsic);
         default:
-          TINT_UNREACHABLE(diagnostics_)
+          TINT_UNREACHABLE(Writer, diagnostics_)
               << "unsupported DecomposeMemoryAccess::Intrinsic storage class:"
               << intrinsic->storage_class;
           return false;
@@ -493,8 +495,9 @@
 
   auto* func = builder_.AST().Functions().Find(ident->symbol());
   if (func == nullptr) {
-    diagnostics_.add_error("Unable to find function: " +
-                           builder_.Symbols().NameFor(ident->symbol()));
+    diagnostics_.add_error(diag::System::Writer,
+                           "Unable to find function: " +
+                               builder_.Symbols().NameFor(ident->symbol()));
     return false;
   }
 
@@ -610,7 +613,7 @@
         case DataType::kVec4I32:
           return cast("asint", load_vec4);
       }
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
           << static_cast<int>(intrinsic->type);
       return false;
@@ -618,7 +621,7 @@
     default:
       break;
   }
-  TINT_UNREACHABLE(diagnostics_)
+  TINT_UNREACHABLE(Writer, diagnostics_)
       << "unsupported DecomposeMemoryAccess::Intrinsic::Op: "
       << static_cast<int>(intrinsic->op);
   return false;
@@ -681,7 +684,7 @@
         case DataType::kVec4I32:
           return load("asint", 4);
       }
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
           << static_cast<int>(intrinsic->type);
       return false;
@@ -733,7 +736,7 @@
         case DataType::kVec4I32:
           return store(4);
       }
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported DecomposeMemoryAccess::Intrinsic::DataType: "
           << static_cast<int>(intrinsic->type);
       return false;
@@ -752,7 +755,7 @@
       return EmitStorageAtomicCall(pre, out, expr, intrinsic->op);
   }
 
-  TINT_UNREACHABLE(diagnostics_)
+  TINT_UNREACHABLE(Writer, diagnostics_)
       << "unsupported DecomposeMemoryAccess::Intrinsic::Op: "
       << static_cast<int>(intrinsic->op);
   return false;
@@ -921,7 +924,7 @@
       break;
 
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
           << static_cast<int>(op);
       return false;
@@ -1077,7 +1080,7 @@
       break;
 
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported atomic intrinsic: " << intrinsic->Type();
       return false;
   }
@@ -1266,6 +1269,7 @@
       break;
     default:
       diagnostics_.add_error(
+          diag::System::Writer,
           "Internal error: unhandled data packing intrinsic");
       return false;
   }
@@ -1337,6 +1341,7 @@
       break;
     default:
       diagnostics_.add_error(
+          diag::System::Writer,
           "Internal error: unhandled data packing intrinsic");
       return false;
   }
@@ -1354,7 +1359,7 @@
   } else if (intrinsic->Type() == sem::IntrinsicType::kStorageBarrier) {
     out << "DeviceMemoryBarrierWithGroupSync()";
   } else {
-    TINT_UNREACHABLE(diagnostics_)
+    TINT_UNREACHABLE(Writer, diagnostics_)
         << "unexpected barrier intrinsic type " << sem::str(intrinsic->Type());
     return false;
   }
@@ -1378,7 +1383,7 @@
 
   auto* texture = arg(Usage::kTexture);
   if (!texture) {
-    TINT_ICE(diagnostics_) << "missing texture argument";
+    TINT_ICE(Writer, diagnostics_) << "missing texture argument";
     return false;
   }
 
@@ -1398,7 +1403,7 @@
         case sem::IntrinsicType::kTextureDimensions:
           switch (texture_type->dim()) {
             case ast::TextureDimension::kNone:
-              TINT_ICE(diagnostics_) << "texture dimension is kNone";
+              TINT_ICE(Writer, diagnostics_) << "texture dimension is kNone";
               return false;
             case ast::TextureDimension::k1d:
               num_dimensions = 1;
@@ -1426,7 +1431,8 @@
         case sem::IntrinsicType::kTextureNumLayers:
           switch (texture_type->dim()) {
             default:
-              TINT_ICE(diagnostics_) << "texture dimension is not arrayed";
+              TINT_ICE(Writer, diagnostics_)
+                  << "texture dimension is not arrayed";
               return false;
             case ast::TextureDimension::k2dArray:
               num_dimensions = is_ms ? 4 : 3;
@@ -1441,7 +1447,7 @@
         case sem::IntrinsicType::kTextureNumLevels:
           switch (texture_type->dim()) {
             default:
-              TINT_ICE(diagnostics_)
+              TINT_ICE(Writer, diagnostics_)
                   << "texture dimension does not support mips";
               return false;
             case ast::TextureDimension::k2d:
@@ -1460,7 +1466,7 @@
         case sem::IntrinsicType::kTextureNumSamples:
           switch (texture_type->dim()) {
             default:
-              TINT_ICE(diagnostics_)
+              TINT_ICE(Writer, diagnostics_)
                   << "texture dimension does not support multisampling";
               return false;
             case ast::TextureDimension::k2d:
@@ -1474,7 +1480,7 @@
           }
           break;
         default:
-          TINT_ICE(diagnostics_) << "unexpected intrinsic";
+          TINT_ICE(Writer, diagnostics_) << "unexpected intrinsic";
           return false;
       }
 
@@ -1496,7 +1502,7 @@
       }
 
       if (num_dimensions > 4) {
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "Texture query intrinsic temporary vector has " << num_dimensions
             << " dimensions";
         return false;
@@ -1533,7 +1539,8 @@
       } else {
         static constexpr char xyzw[] = {'x', 'y', 'z', 'w'};
         if (num_dimensions < 0 || num_dimensions > 4) {
-          TINT_ICE(diagnostics_) << "vector dimensions are " << num_dimensions;
+          TINT_ICE(Writer, diagnostics_)
+              << "vector dimensions are " << num_dimensions;
           return false;
         }
         for (int i = 0; i < num_dimensions; i++) {
@@ -1593,8 +1600,9 @@
       break;
     default:
       diagnostics_.add_error(
+          diag::System::Writer,
           "Internal compiler error: Unhandled texture intrinsic '" +
-          std::string(intrinsic->str()) + "'");
+              std::string(intrinsic->str()) + "'");
       return false;
   }
 
@@ -1606,7 +1614,7 @@
 
   auto* param_coords = arg(Usage::kCoords);
   if (!param_coords) {
-    TINT_ICE(diagnostics_) << "missing coords argument";
+    TINT_ICE(Writer, diagnostics_) << "missing coords argument";
     return false;
   }
 
@@ -1674,9 +1682,10 @@
       }
     }
     if (wgsl_ret_width > hlsl_ret_width) {
-      TINT_ICE(diagnostics_) << "WGSL return width (" << wgsl_ret_width
-                             << ") is wider than HLSL return width ("
-                             << hlsl_ret_width << ") for " << intrinsic->Type();
+      TINT_ICE(Writer, diagnostics_)
+          << "WGSL return width (" << wgsl_ret_width
+          << ") is wider than HLSL return width (" << hlsl_ret_width << ") for "
+          << intrinsic->Type();
       return false;
     }
   }
@@ -1766,8 +1775,9 @@
     case sem::IntrinsicType::kSmoothStep:
       return "smoothstep";
     default:
-      diagnostics_.add_error("Unknown builtin method: " +
-                             std::string(intrinsic->str()));
+      diagnostics_.add_error(
+          diag::System::Writer,
+          "Unknown builtin method: " + std::string(intrinsic->str()));
   }
 
   return "";
@@ -1936,7 +1946,8 @@
     return EmitUnaryOp(pre, out, u);
   }
 
-  diagnostics_.add_error("unknown expression type: " + builder_.str(expr));
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown expression type: " + builder_.str(expr));
   return false;
 }
 
@@ -2084,7 +2095,8 @@
       break;
   }
 
-  TINT_ICE(diagnostics_) << "unhandled storage class " << sem->StorageClass();
+  TINT_ICE(Writer, diagnostics_)
+      << "unhandled storage class " << sem->StorageClass();
   return false;
 }
 
@@ -2099,7 +2111,7 @@
   auto* str = type->As<sem::Struct>();
   if (!str) {
     // https://www.w3.org/TR/WGSL/#module-scope-variables
-    TINT_ICE(diagnostics_)
+    TINT_ICE(Writer, diagnostics_)
         << "variables with uniform storage must be structure";
   }
 
@@ -2285,7 +2297,7 @@
       if (wgsize[i].overridable_const) {
         auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
         if (!sem_const->IsPipelineConstant()) {
-          TINT_ICE(builder_.Diagnostics())
+          TINT_ICE(Writer, builder_.Diagnostics())
               << "expected a pipeline-overridable constant";
         }
         out << kSpecConstantPrefix << sem_const->ConstantId();
@@ -2310,7 +2322,8 @@
     if (!type->Is<sem::Struct>()) {
       // ICE likely indicates that the CanonicalizeEntryPointIO transform was
       // not run, or a builtin parameter was added after it was run.
-      TINT_ICE(diagnostics_) << "Unsupported non-struct entry point parameter";
+      TINT_ICE(Writer, diagnostics_)
+          << "Unsupported non-struct entry point parameter";
     }
 
     if (!first) {
@@ -2359,7 +2372,7 @@
   } else if (auto* ul = lit->As<ast::UintLiteral>()) {
     out << ul->value() << "u";
   } else {
-    diagnostics_.add_error("unknown literal type");
+    diagnostics_.add_error(diag::System::Writer, "unknown literal type");
     return false;
   }
   return true;
@@ -2435,8 +2448,9 @@
     }
     out << "}";
   } else {
-    diagnostics_.add_error("Invalid type for zero emission: " +
-                           type->type_name());
+    diagnostics_.add_error(
+        diag::System::Writer,
+        "Invalid type for zero emission: " + type->type_name());
     return false;
   }
   return true;
@@ -2556,7 +2570,8 @@
     return EmitVariable(out, v->variable());
   }
 
-  diagnostics_.add_error("unknown statement type: " + builder_.str(stmt));
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown statement type: " + builder_.str(stmt));
   return false;
 }
 
@@ -2604,7 +2619,7 @@
       auto* str = type->As<sem::Struct>();
       if (!str) {
         // https://www.w3.org/TR/WGSL/#module-scope-variables
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "variables with uniform storage must be structure";
       }
       auto array_length = (str->Size() + 15) / 16;
@@ -2623,7 +2638,7 @@
     std::vector<uint32_t> sizes;
     while (auto* arr = base_type->As<sem::Array>()) {
       if (arr->IsRuntimeSized()) {
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "Runtime arrays may only exist in storage buffers, which should "
                "have been transformed into a ByteAddressBuffer";
         return false;
@@ -2662,7 +2677,7 @@
     // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-per-component-math#matrix-ordering
     out << mat->columns() << "x" << mat->rows();
   } else if (type->Is<sem::Pointer>()) {
-    TINT_ICE(diagnostics_)
+    TINT_ICE(Writer, diagnostics_)
         << "Attempting to emit pointer type. These should have been removed "
            "with the InlinePointerLets transform";
     return false;
@@ -2704,7 +2719,7 @@
         out << "CubeArray";
         break;
       default:
-        TINT_UNREACHABLE(diagnostics_)
+        TINT_UNREACHABLE(Writer, diagnostics_)
             << "unexpected TextureDimension " << tex->dim();
         return false;
     }
@@ -2712,8 +2727,9 @@
     if (storage) {
       auto* component = image_format_to_rwtexture_type(storage->image_format());
       if (component == nullptr) {
-        TINT_ICE(diagnostics_) << "Unsupported StorageTexture ImageFormat: "
-                               << static_cast<int>(storage->image_format());
+        TINT_ICE(Writer, diagnostics_)
+            << "Unsupported StorageTexture ImageFormat: "
+            << static_cast<int>(storage->image_format());
         return false;
       }
       out << "<" << component << ">";
@@ -2727,7 +2743,8 @@
       } else if (subtype->Is<sem::U32>()) {
         out << "uint4";
       } else {
-        TINT_ICE(diagnostics_) << "Unsupported multisampled texture type";
+        TINT_ICE(Writer, diagnostics_)
+            << "Unsupported multisampled texture type";
         return false;
       }
       out << ">";
@@ -2758,7 +2775,7 @@
   } else if (type->Is<sem::Void>()) {
     out << "void";
   } else {
-    diagnostics_.add_error("unknown type in EmitType");
+    diagnostics_.add_error(diag::System::Writer, "unknown type in EmitType");
     return false;
   }
 
@@ -2814,7 +2831,8 @@
       if (auto* location = deco->As<ast::LocationDecoration>()) {
         auto& pipeline_stage_uses = str->PipelineStageUses();
         if (pipeline_stage_uses.size() != 1) {
-          TINT_ICE(diagnostics_) << "invalid entry point IO struct uses";
+          TINT_ICE(Writer, diagnostics_)
+              << "invalid entry point IO struct uses";
         }
 
         if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
@@ -2829,12 +2847,13 @@
                        sem::PipelineStageUsage::kFragmentOutput)) {
           out << " : SV_Target" + std::to_string(location->value());
         } else {
-          TINT_ICE(diagnostics_) << "invalid use of location decoration";
+          TINT_ICE(Writer, diagnostics_)
+              << "invalid use of location decoration";
         }
       } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
         auto attr = builtin_to_attribute(builtin->value());
         if (attr.empty()) {
-          diagnostics_.add_error("unsupported builtin");
+          diagnostics_.add_error(diag::System::Writer, "unsupported builtin");
           return false;
         }
         out << " : " << attr;
@@ -2888,7 +2907,8 @@
 
   // TODO(dsinclair): Handle variable decorations
   if (!var->decorations().empty()) {
-    diagnostics_.add_error("Variable decorations are not handled yet");
+    diagnostics_.add_error(diag::System::Writer,
+                           "Variable decorations are not handled yet");
     return false;
   }
 
@@ -2925,12 +2945,13 @@
 
   for (auto* d : var->decorations()) {
     if (!d->Is<ast::OverrideDecoration>()) {
-      diagnostics_.add_error("Decorated const values not valid");
+      diagnostics_.add_error(diag::System::Writer,
+                             "Decorated const values not valid");
       return false;
     }
   }
   if (!var->is_const()) {
-    diagnostics_.add_error("Expected a const value");
+    diagnostics_.add_error(diag::System::Writer, "Expected a const value");
     return false;
   }
 
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 6e74487..5a89d11 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -101,7 +101,7 @@
         case ast::StorageClass::kPrivate:
         case ast::StorageClass::kWorkgroup:
           // These are pushed into the entry point by the sanitizer.
-          TINT_ICE(diagnostics_)
+          TINT_ICE(Writer, diagnostics_)
               << "module-scope variables in the private/workgroup storage "
                  "class should have been handled by the MSL sanitizer";
           break;
@@ -134,7 +134,8 @@
       return false;
     }
   } else {
-    diagnostics_.add_error("unknown alias type: " + ty->type_name());
+    diagnostics_.add_error(diag::System::Writer,
+                           "unknown alias type: " + ty->type_name());
     return false;
   }
 
@@ -269,7 +270,8 @@
       out_ << "%";
       break;
     case ast::BinaryOp::kNone:
-      diagnostics_.add_error("missing binary operation type");
+      diagnostics_.add_error(diag::System::Writer,
+                             "missing binary operation type");
       return false;
   }
   out_ << " ";
@@ -297,8 +299,9 @@
 
   auto* func = program_->AST().Functions().Find(ident->symbol());
   if (func == nullptr) {
-    diagnostics_.add_error("Unable to find function: " +
-                           program_->Symbols().NameFor(ident->symbol()));
+    diagnostics_.add_error(diag::System::Writer,
+                           "Unable to find function: " +
+                               program_->Symbols().NameFor(ident->symbol()));
     return false;
   }
 
@@ -425,7 +428,7 @@
 
   auto* texture = arg(Usage::kTexture);
   if (!texture) {
-    TINT_ICE(diagnostics_) << "missing texture arg";
+    TINT_ICE(Writer, diagnostics_) << "missing texture arg";
     return false;
   }
 
@@ -457,7 +460,8 @@
       std::vector<const char*> dims;
       switch (texture_type->dim()) {
         case ast::TextureDimension::kNone:
-          diagnostics_.add_error("texture dimension is kNone");
+          diagnostics_.add_error(diag::System::Writer,
+                                 "texture dimension is kNone");
           return false;
         case ast::TextureDimension::k1d:
           dims = {"width"};
@@ -557,7 +561,7 @@
       out_ << ".write(";
       break;
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "Unhandled texture intrinsic '" << intrinsic->str() << "'";
       return false;
   }
@@ -595,7 +599,8 @@
             out_ << "uint3(";
             break;
           default:
-            TINT_ICE(diagnostics_) << "unhandled texture dimensionality";
+            TINT_ICE(Writer, diagnostics_)
+                << "unhandled texture dimensionality";
             break;
         }
       }
@@ -653,7 +658,7 @@
       default: {
         std::stringstream err;
         err << "MSL does not support gradients for " << dim << " textures";
-        diagnostics_.add_error(err.str());
+        diagnostics_.add_error(diag::System::Writer, err.str());
         return false;
       }
     }
@@ -813,8 +818,9 @@
       out += "unpack_unorm2x16_to_float";
       break;
     default:
-      diagnostics_.add_error("Unknown import method: " +
-                             std::string(intrinsic->str()));
+      diagnostics_.add_error(
+          diag::System::Writer,
+          "Unknown import method: " + std::string(intrinsic->str()));
       return "";
   }
   return out;
@@ -949,8 +955,9 @@
   } else if (type->As<sem::Struct>()) {
     out_ << "{}";
   } else {
-    diagnostics_.add_error("Invalid type for zero emission: " +
-                           type->type_name());
+    diagnostics_.add_error(
+        diag::System::Writer,
+        "Invalid type for zero emission: " + type->type_name());
     return false;
   }
   return true;
@@ -971,7 +978,7 @@
   } else if (auto* ul = lit->As<ast::UintLiteral>()) {
     out_ << ul->value() << "u";
   } else {
-    diagnostics_.add_error("unknown literal type");
+    diagnostics_.add_error(diag::System::Writer, "unknown literal type");
     return false;
   }
   return true;
@@ -1003,7 +1010,8 @@
     return EmitUnaryOp(u);
   }
 
-  diagnostics_.add_error("unknown expression type: " + program_->str(expr));
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown expression type: " + program_->str(expr));
   return false;
 }
 
@@ -1159,7 +1167,7 @@
       auto* binding =
           ast::GetDecoration<ast::BindingDecoration>(var->decorations());
       if (binding == nullptr) {
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "missing binding attribute for entry point parameter";
         return false;
       }
@@ -1168,7 +1176,8 @@
       } else if (var->type()->Is<ast::Texture>()) {
         out_ << " [[texture(" << binding->value() << ")]]";
       } else {
-        TINT_ICE(diagnostics_) << "invalid handle type entry point parameter";
+        TINT_ICE(Writer, diagnostics_)
+            << "invalid handle type entry point parameter";
         return false;
       }
     } else {
@@ -1184,13 +1193,13 @@
 
         auto attr = builtin_to_attribute(builtin->value());
         if (attr.empty()) {
-          diagnostics_.add_error("unknown builtin");
+          diagnostics_.add_error(diag::System::Writer, "unknown builtin");
           return false;
         }
         out_ << " [[" << attr << "]]";
       }
       if (!builtin_found) {
-        TINT_ICE(diagnostics_) << "Unsupported entry point parameter";
+        TINT_ICE(Writer, diagnostics_) << "Unsupported entry point parameter";
       }
     }
   }
@@ -1208,8 +1217,9 @@
     auto* binding = data.second.binding;
     if (binding == nullptr) {
       diagnostics_.add_error(
+          diag::System::Writer,
           "unable to find binding information for uniform: " +
-          program_->Symbols().NameFor(var->Declaration()->symbol()));
+              program_->Symbols().NameFor(var->Declaration()->symbol()));
       return false;
     }
     // auto* set = data.second.set;
@@ -1472,7 +1482,8 @@
     return EmitVariable(var);
   }
 
-  diagnostics_.add_error("unknown statement type: " + program_->str(stmt));
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown statement type: " + program_->str(stmt));
   return false;
 }
 
@@ -1549,7 +1560,7 @@
         out_ << "constant ";
         break;
       default:
-        TINT_ICE(diagnostics_) << "unhandled storage class for pointer";
+        TINT_ICE(Writer, diagnostics_) << "unhandled storage class for pointer";
     }
     if (ptr->StoreType()->Is<sem::Array>()) {
       std::string inner = "(*" + name + ")";
@@ -1595,7 +1606,8 @@
         out_ << "cube_array";
         break;
       default:
-        diagnostics_.add_error("Invalid texture dimensions");
+        diagnostics_.add_error(diag::System::Writer,
+                               "Invalid texture dimensions");
         return false;
     }
     if (tex->Is<sem::MultisampledTexture>()) {
@@ -1615,7 +1627,8 @@
       } else if (storage->access() == ast::Access::kWrite) {
         out_ << ", access::write";
       } else {
-        diagnostics_.add_error("Invalid access control for storage texture");
+        diagnostics_.add_error(diag::System::Writer,
+                               "Invalid access control for storage texture");
         return false;
       }
     } else if (auto* ms = tex->As<sem::MultisampledTexture>()) {
@@ -1629,7 +1642,7 @@
       }
       out_ << ", access::sample";
     } else {
-      diagnostics_.add_error("invalid texture type");
+      diagnostics_.add_error(diag::System::Writer, "invalid texture type");
       return false;
     }
     out_ << ">";
@@ -1644,7 +1657,8 @@
   } else if (type->Is<sem::Void>()) {
     out_ << "void";
   } else {
-    diagnostics_.add_error("unknown type in EmitType: " + type->type_name());
+    diagnostics_.add_error(diag::System::Writer,
+                           "unknown type in EmitType: " + type->type_name());
     return false;
   }
 
@@ -1700,7 +1714,7 @@
     if (is_host_shareable) {
       if (wgsl_offset < msl_offset) {
         // Unimplementable layout
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "Structure member WGSL offset (" << wgsl_offset
             << ") is behind MSL offset (" << msl_offset << ")";
         return false;
@@ -1737,14 +1751,15 @@
       if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
         auto attr = builtin_to_attribute(builtin->value());
         if (attr.empty()) {
-          diagnostics_.add_error("unknown builtin");
+          diagnostics_.add_error(diag::System::Writer, "unknown builtin");
           return false;
         }
         out_ << " [[" << attr << "]]";
       } else if (auto* loc = deco->As<ast::LocationDecoration>()) {
         auto& pipeline_stage_uses = str->PipelineStageUses();
         if (pipeline_stage_uses.size() != 1) {
-          TINT_ICE(diagnostics_) << "invalid entry point IO struct uses";
+          TINT_ICE(Writer, diagnostics_)
+              << "invalid entry point IO struct uses";
         }
 
         if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
@@ -1759,7 +1774,8 @@
                        sem::PipelineStageUsage::kFragmentOutput)) {
           out_ << " [[color(" + std::to_string(loc->value()) + ")]]";
         } else {
-          TINT_ICE(diagnostics_) << "invalid use of location decoration";
+          TINT_ICE(Writer, diagnostics_)
+              << "invalid use of location decoration";
         }
       }
     }
@@ -1770,7 +1786,7 @@
       // Calculate new MSL offset
       auto size_align = MslPackedTypeSizeAndAlign(ty);
       if (msl_offset % size_align.align) {
-        TINT_ICE(diagnostics_)
+        TINT_ICE(Writer, diagnostics_)
             << "Misaligned MSL structure member "
             << ty->FriendlyName(program_->Symbols()) << " " << name;
         return false;
@@ -1828,7 +1844,7 @@
 
   for (auto* deco : decl->decorations()) {
     if (!deco->Is<ast::InternalDecoration>()) {
-      TINT_ICE(diagnostics_) << "unexpected variable decoration";
+      TINT_ICE(Writer, diagnostics_) << "unexpected variable decoration";
       return false;
     }
   }
@@ -1845,7 +1861,7 @@
       out_ << "threadgroup ";
       break;
     default:
-      TINT_ICE(diagnostics_) << "unhandled variable storage class";
+      TINT_ICE(Writer, diagnostics_) << "unhandled variable storage class";
       return false;
   }
 
@@ -1886,12 +1902,13 @@
 
   for (auto* d : var->decorations()) {
     if (!d->Is<ast::OverrideDecoration>()) {
-      diagnostics_.add_error("Decorated const values not valid");
+      diagnostics_.add_error(diag::System::Writer,
+                             "Decorated const values not valid");
       return false;
     }
   }
   if (!var->is_const()) {
-    diagnostics_.add_error("Expected a const value");
+    diagnostics_.add_error(diag::System::Writer, "Expected a const value");
     return false;
   }
 
@@ -1963,8 +1980,9 @@
 
   if (auto* arr = ty->As<sem::Array>()) {
     if (!arr->IsStrideImplicit()) {
-      TINT_ICE(diagnostics_) << "arrays with explicit strides should have "
-                                "removed with the PadArrayElements transform";
+      TINT_ICE(Writer, diagnostics_)
+          << "arrays with explicit strides should have "
+             "removed with the PadArrayElements transform";
       return {};
     }
     auto num_els = std::max<uint32_t>(arr->Count(), 1);
@@ -1977,7 +1995,8 @@
     return SizeAndAlign{str->Size(), str->Align()};
   }
 
-  TINT_UNREACHABLE(diagnostics_) << "Unhandled type " << ty->TypeInfo().name;
+  TINT_UNREACHABLE(Writer, diagnostics_)
+      << "Unhandled type " << ty->TypeInfo().name;
   return {};
 }
 
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 463e822..450afb2 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -453,7 +453,7 @@
       if (has_overridable_workgroup_size_) {
         // Only one stage can have a pipeline-overridable workgroup size.
         // TODO(crbug.com/tint/810): Use LocalSizeId to handle this scenario.
-        TINT_ICE(builder_.Diagnostics())
+        TINT_ICE(Writer, builder_.Diagnostics())
             << "multiple stages using pipeline-overridable workgroup sizes";
       }
       has_overridable_workgroup_size_ = true;
@@ -477,7 +477,7 @@
           // Make the constant specializable.
           auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
           if (!sem_const->IsPipelineConstant()) {
-            TINT_ICE(builder_.Diagnostics())
+            TINT_ICE(Writer, builder_.Diagnostics())
                 << "expected a pipeline-overridable constant";
           }
           constant.is_spec_op = true;
@@ -885,7 +885,8 @@
   if (auto* scalar = expr->idx_expr()->As<ast::ScalarConstructorExpression>()) {
     auto* literal = scalar->literal()->As<ast::IntLiteral>();
     if (!literal) {
-      TINT_ICE(builder_.Diagnostics()) << "bad literal in array accessor";
+      TINT_ICE(Writer, builder_.Diagnostics())
+          << "bad literal in array accessor";
       return false;
     }
 
@@ -917,7 +918,8 @@
     return true;
   }
 
-  TINT_ICE(builder_.Diagnostics()) << "unsupported array accessor expression";
+  TINT_ICE(Writer, builder_.Diagnostics())
+      << "unsupported array accessor expression";
   return false;
 }
 
@@ -1042,7 +1044,7 @@
     return true;
   }
 
-  TINT_ICE(builder_.Diagnostics())
+  TINT_ICE(Writer, builder_.Diagnostics())
       << "unhandled member index type: " << expr_sem->TypeInfo().name;
   return false;
 }
@@ -1050,7 +1052,7 @@
 uint32_t Builder::GenerateAccessorExpression(ast::Expression* expr) {
   if (!expr->IsAnyOf<ast::ArrayAccessorExpression,
                      ast::MemberAccessorExpression>()) {
-    TINT_ICE(builder_.Diagnostics()) << "expression is not an accessor";
+    TINT_ICE(Writer, builder_.Diagnostics()) << "expression is not an accessor";
     return 0;
   }
 
@@ -1466,7 +1468,7 @@
   // This should not happen as we rely on constant folding to obviate
   // casts/conversions for module-scope variables
   if (is_global_init) {
-    TINT_ICE(builder_.Diagnostics())
+    TINT_ICE(Writer, builder_.Diagnostics())
         << "Module-level conversions are not supported. Conversions should "
            "have already been constant-folded by the FoldConstants transform.";
     return 0;
@@ -1545,7 +1547,7 @@
     return result_id;
 
   } else {
-    TINT_ICE(builder_.Diagnostics()) << "Invalid from_type";
+    TINT_ICE(Writer, builder_.Diagnostics()) << "Invalid from_type";
   }
 
   if (op == spv::Op::OpNop) {
@@ -2506,7 +2508,7 @@
   auto gen_arg = [&](Usage usage) {
     auto* argument = arg(usage);
     if (!argument) {
-      TINT_ICE(builder_.Diagnostics())
+      TINT_ICE(Writer, builder_.Diagnostics())
           << "missing argument " << static_cast<int>(usage);
     }
     return gen(argument);
@@ -2514,7 +2516,7 @@
 
   auto* texture = arg(Usage::kTexture);
   if (!texture) {
-    TINT_ICE(builder_.Diagnostics()) << "missing texture argument";
+    TINT_ICE(Writer, builder_.Diagnostics()) << "missing texture argument";
   }
 
   auto* texture_type = TypeOf(texture)->UnwrapRef()->As<sem::Texture>();
@@ -2845,7 +2847,7 @@
       break;
     }
     default:
-      TINT_UNREACHABLE(builder_.Diagnostics());
+      TINT_UNREACHABLE(Writer, builder_.Diagnostics());
       return false;
   }
 
@@ -2942,7 +2944,7 @@
           ScalarConstant::U32(static_cast<uint32_t>(spv::Scope::Device)));
       break;
     default:
-      TINT_UNREACHABLE(builder_.Diagnostics())
+      TINT_UNREACHABLE(Writer, builder_.Diagnostics())
           << "unhandled atomic storage class " << storage_class;
       return false;
   }
@@ -3115,7 +3117,7 @@
         zero = GenerateConstantIfNeeded(ScalarConstant::U32(0u));
         one = GenerateConstantIfNeeded(ScalarConstant::U32(1u));
       } else {
-        TINT_UNREACHABLE(builder_.Diagnostics())
+        TINT_UNREACHABLE(Writer, builder_.Diagnostics())
             << "unsupported atomic type " << value_sem_type->TypeInfo().name;
       }
       if (zero == 0 || one == 0) {
@@ -3144,7 +3146,7 @@
                                 });
     }
     default:
-      TINT_UNREACHABLE(builder_.Diagnostics())
+      TINT_UNREACHABLE(Writer, builder_.Diagnostics())
           << "unhandled atomic intrinsic " << intrinsic->Type();
       return false;
   }
@@ -3374,7 +3376,7 @@
     if (LastIsFallthrough(item->body())) {
       if (i == (body.size() - 1)) {
         // This case is caught by Resolver validation
-        TINT_UNREACHABLE(builder_.Diagnostics());
+        TINT_UNREACHABLE(Writer, builder_.Diagnostics());
         return false;
       }
       if (!push_function_inst(spv::Op::OpBranch,
@@ -3938,7 +3940,8 @@
       } else if (storage == ast::StorageClass::kOutput) {
         return SpvBuiltInPosition;
       } else {
-        TINT_ICE(builder_.Diagnostics()) << "invalid storage class for builtin";
+        TINT_ICE(Writer, builder_.Diagnostics())
+            << "invalid storage class for builtin";
         break;
       }
     case ast::Builtin::kVertexIndex:
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index 79932c3..7c9a6ac 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -192,7 +192,7 @@
   /// @param operands the variable operands
   void push_function_var(const OperandList& operands) {
     if (functions_.empty()) {
-      TINT_ICE(builder_.Diagnostics())
+      TINT_ICE(Writer, builder_.Diagnostics())
           << "push_function_var() called without a function";
     }
     functions_.back().push_var(operands);
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 94d7b3a..82875c4 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -80,7 +80,7 @@
         return false;
       }
     } else {
-      TINT_UNREACHABLE(diagnostics_);
+      TINT_UNREACHABLE(Writer, diagnostics_);
       return false;
     }
 
@@ -106,8 +106,9 @@
       return false;
     }
   } else {
-    diagnostics_.add_error("unknown declared type: " +
-                           std::string(ty->TypeInfo().name));
+    diagnostics_.add_error(
+        diag::System::Writer,
+        "unknown declared type: " + std::string(ty->TypeInfo().name));
     return false;
   }
   return true;
@@ -139,7 +140,7 @@
     return EmitUnaryOp(u);
   }
 
-  diagnostics_.add_error("unknown expression type");
+  diagnostics_.add_error(diag::System::Writer, "unknown expression type");
   return false;
 }
 
@@ -273,7 +274,7 @@
   } else if (auto* ul = lit->As<ast::UintLiteral>()) {
     out_ << ul->value() << "u";
   } else {
-    diagnostics_.add_error("unknown literal type");
+    diagnostics_.add_error(diag::System::Writer, "unknown literal type");
     return false;
   }
   return true;
@@ -349,7 +350,7 @@
 bool GeneratorImpl::EmitImageFormat(const ast::ImageFormat fmt) {
   switch (fmt) {
     case ast::ImageFormat::kNone:
-      diagnostics_.add_error("unknown image format");
+      diagnostics_.add_error(diag::System::Writer, "unknown image format");
       return false;
     default:
       out_ << fmt;
@@ -371,7 +372,7 @@
     default:
       break;
   }
-  diagnostics_.add_error("unknown access");
+  diagnostics_.add_error(diag::System::Writer, "unknown access");
   return false;
 }
 
@@ -435,7 +436,7 @@
     } else if (texture->Is<ast::StorageTexture>()) {
       out_ << "storage_";
     } else {
-      diagnostics_.add_error("unknown texture type");
+      diagnostics_.add_error(diag::System::Writer, "unknown texture type");
       return false;
     }
 
@@ -459,7 +460,8 @@
         out_ << "cube_array";
         break;
       default:
-        diagnostics_.add_error("unknown texture dimension");
+        diagnostics_.add_error(diag::System::Writer,
+                               "unknown texture dimension");
         return false;
     }
 
@@ -500,7 +502,8 @@
   } else if (auto* tn = ty->As<ast::TypeName>()) {
     out_ << program_->Symbols().NameFor(tn->name());
   } else {
-    diagnostics_.add_error("unknown type in EmitType: " + ty->type_name());
+    diagnostics_.add_error(diag::System::Writer,
+                           "unknown type in EmitType: " + ty->type_name());
     return false;
   }
   return true;
@@ -647,7 +650,8 @@
               return false;
             }
           } else {
-            TINT_ICE(diagnostics_) << "Unsupported workgroup_size expression";
+            TINT_ICE(Writer, diagnostics_)
+                << "Unsupported workgroup_size expression";
           }
         }
       }
@@ -676,7 +680,7 @@
     } else if (auto* internal = deco->As<ast::InternalDecoration>()) {
       out_ << "internal(" << internal->InternalName() << ")";
     } else {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Writer, diagnostics_)
           << "Unsupported decoration '" << deco->TypeInfo().name << "'";
       return false;
     }
@@ -750,7 +754,8 @@
       out_ << "%";
       break;
     case ast::BinaryOp::kNone:
-      diagnostics_.add_error("missing binary operation type");
+      diagnostics_.add_error(diag::System::Writer,
+                             "missing binary operation type");
       return false;
   }
   out_ << " ";
@@ -870,7 +875,8 @@
     return EmitVariable(v->variable());
   }
 
-  diagnostics_.add_error("unknown statement type: " + program_->str(stmt));
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown statement type: " + program_->str(stmt));
   return false;
 }
 
@@ -1059,7 +1065,8 @@
       return ident;
     }
   }
-  diagnostics_.add_error("Unable to generate a unique WGSL identifier");
+  diagnostics_.add_error(diag::System::Writer,
+                         "Unable to generate a unique WGSL identifier");
   return "<invalid-ident>";
 }