diag: Add System enumerator to all diagnostics

Describes what Tint system raised the diagnostic.

Use this information in the fuzzers to distinguish between expected and unexpected failure cases in the Transform fuzzer tests.

Fixed: chromium:1206407
Fixed: chromium:1207154
Change-Id: I3b807acafe384a2fc363d2a4165a29693450b3cf
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55254
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>

fuzzers/tint_common_fuzzer.cc

@@ -35,12 +35,13 @@ namespace {
   __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 @@ int CommonFuzzer::Run(const uint8_t* data, size_t size) {
   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);

src/ast/alias.cc

@@ -28,7 +28,7 @@ Alias::Alias(ProgramID program_id,
     : 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;

src/ast/array.cc

@@ -38,7 +38,7 @@ Array::Array(Array&&) = default;
 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()) {

src/ast/array_accessor_expression.cc

@@ -26,10 +26,10 @@ ArrayAccessorExpression::ArrayAccessorExpression(ProgramID program_id,
                                                  Expression* array,
                                                  Expression* idx_expr)
     : Base(program_id, source), array_(array), idx_expr_(idx_expr) {
-  TINT_ASSERT(array_);
+  TINT_ASSERT(AST, array_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(array_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, array_, program_id);
-  TINT_ASSERT(idx_expr_);
+  TINT_ASSERT(AST, idx_expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(idx_expr_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, idx_expr_, program_id);
 }
 
 ArrayAccessorExpression::ArrayAccessorExpression(ArrayAccessorExpression&&) =

src/ast/assignment_statement.cc

@@ -26,10 +26,10 @@ AssignmentStatement::AssignmentStatement(ProgramID program_id,
                                          Expression* lhs,
                                          Expression* rhs)
     : Base(program_id, source), lhs_(lhs), rhs_(rhs) {
-  TINT_ASSERT(lhs_);
+  TINT_ASSERT(AST, lhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(lhs_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs_, program_id);
-  TINT_ASSERT(rhs_);
+  TINT_ASSERT(AST, rhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(rhs_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs_, program_id);
 }
 
 AssignmentStatement::AssignmentStatement(AssignmentStatement&&) = default;

src/ast/binary_expression.cc

@@ -27,11 +27,11 @@ BinaryExpression::BinaryExpression(ProgramID program_id,
                                    Expression* lhs,
                                    Expression* rhs)
     : Base(program_id, source), op_(op), lhs_(lhs), rhs_(rhs) {
-  TINT_ASSERT(lhs_);
+  TINT_ASSERT(AST, lhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(lhs_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs_, program_id);
-  TINT_ASSERT(rhs_);
+  TINT_ASSERT(AST, rhs_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(rhs_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs_, program_id);
-  TINT_ASSERT(op_ != BinaryOp::kNone);
+  TINT_ASSERT(AST, op_ != BinaryOp::kNone);
 }
 
 BinaryExpression::BinaryExpression(BinaryExpression&&) = default;

src/ast/bitcast_expression.cc

@@ -26,9 +26,9 @@ BitcastExpression::BitcastExpression(ProgramID program_id,
                                      ast::Type* type,
                                      Expression* expr)
     : Base(program_id, source), type_(type), expr_(expr) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
-  TINT_ASSERT(expr_);
+  TINT_ASSERT(AST, expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(expr, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, expr, program_id);
 }
 
 BitcastExpression::BitcastExpression(BitcastExpression&&) = default;

src/ast/block_statement.cc

@@ -26,8 +26,8 @@ BlockStatement::BlockStatement(ProgramID program_id,
                                const StatementList& statements)
     : Base(program_id, source), statements_(std::move(statements)) {
   for (auto* stmt : *this) {
-    TINT_ASSERT(stmt);
+    TINT_ASSERT(AST, stmt);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(stmt, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, stmt, program_id);
   }
 }
 

src/ast/call_expression.cc

@@ -26,11 +26,11 @@ CallExpression::CallExpression(ProgramID program_id,
                                IdentifierExpression* func,
                                ExpressionList params)
     : Base(program_id, source), func_(func), params_(params) {
-  TINT_ASSERT(func_);
+  TINT_ASSERT(AST, func_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(func_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, func_, program_id);
   for (auto* param : params_) {
-    TINT_ASSERT(param);
+    TINT_ASSERT(AST, param);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(param, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, param, program_id);
   }
 }
 

src/ast/call_statement.cc

@@ -25,8 +25,8 @@ CallStatement::CallStatement(ProgramID program_id,
                              const Source& source,
                              CallExpression* call)
     : Base(program_id, source), call_(call) {
-  TINT_ASSERT(call_);
+  TINT_ASSERT(AST, call_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(call_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, call_, program_id);
 }
 
 CallStatement::CallStatement(CallStatement&&) = default;

src/ast/case_statement.cc

@@ -26,11 +26,11 @@ CaseStatement::CaseStatement(ProgramID program_id,
                              CaseSelectorList selectors,
                              BlockStatement* body)
     : Base(program_id, source), selectors_(selectors), body_(body) {
-  TINT_ASSERT(body_);
+  TINT_ASSERT(AST, body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
   for (auto* selector : selectors) {
-    TINT_ASSERT(selector);
+    TINT_ASSERT(AST, selector);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(selector, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, selector, program_id);
   }
 }
 

src/ast/depth_texture.cc

@@ -33,7 +33,7 @@ DepthTexture::DepthTexture(ProgramID program_id,
                            const Source& source,
                            TextureDimension dim)
     : Base(program_id, source, dim) {
-  TINT_ASSERT(IsValidDepthDimension(dim));
+  TINT_ASSERT(AST, IsValidDepthDimension(dim));
 }
 
 DepthTexture::DepthTexture(DepthTexture&&) = default;

src/ast/else_statement.cc

@@ -26,9 +26,9 @@ ElseStatement::ElseStatement(ProgramID program_id,
                              Expression* condition,
                              BlockStatement* body)
     : Base(program_id, source), condition_(condition), body_(body) {
-  TINT_ASSERT(body_);
+  TINT_ASSERT(AST, body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
 }
 
 ElseStatement::ElseStatement(ElseStatement&&) = default;

src/ast/function.cc

@@ -38,19 +38,19 @@ Function::Function(ProgramID program_id,
       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(AST, symbol_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body, 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(AST, param && param->is_const());
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(param, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, param, program_id);
   }
-  TINT_ASSERT(symbol_.IsValid());
+  TINT_ASSERT(AST, symbol_.IsValid());
-  TINT_ASSERT(return_type_);
+  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);
   }
 }
 

src/ast/identifier_expression.cc

@@ -25,8 +25,8 @@ IdentifierExpression::IdentifierExpression(ProgramID program_id,
                                            const Source& source,
                                            Symbol sym)
     : Base(program_id, source), sym_(sym) {
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(sym_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, sym_, program_id);
-  TINT_ASSERT(sym_.IsValid());
+  TINT_ASSERT(AST, sym_.IsValid());
 }
 
 IdentifierExpression::IdentifierExpression(IdentifierExpression&&) = default;

src/ast/if_statement.cc

@@ -30,13 +30,13 @@ IfStatement::IfStatement(ProgramID program_id,
       condition_(condition),
       body_(body),
       else_statements_(std::move(else_stmts)) {
-  TINT_ASSERT(condition_);
+  TINT_ASSERT(AST, condition_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
-  TINT_ASSERT(body_);
+  TINT_ASSERT(AST, body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
   for (auto* el : else_statements_) {
-    TINT_ASSERT(el);
+    TINT_ASSERT(AST, el);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(el, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, el, program_id);
   }
 }
 

src/ast/intrinsic_texture_helper_test.cc

@@ -142,7 +142,7 @@ ast::Type* TextureOverloadCase::buildResultVectorComponentType(
       return b->ty.i32();
   }
 
-  TINT_UNREACHABLE(b->Diagnostics());
+  TINT_UNREACHABLE(AST, b->Diagnostics());
   return {};
 }
 
@@ -176,7 +176,7 @@ ast::Variable* TextureOverloadCase::buildTextureVariable(
     }
   }
 
-  TINT_UNREACHABLE(b->Diagnostics());
+  TINT_UNREACHABLE(AST, b->Diagnostics());
   return nullptr;
 }
 

src/ast/loop_statement.cc

@@ -26,9 +26,9 @@ LoopStatement::LoopStatement(ProgramID program_id,
                              BlockStatement* body,
                              BlockStatement* continuing)
     : Base(program_id, source), body_(body), continuing_(continuing) {
-  TINT_ASSERT(body_);
+  TINT_ASSERT(AST, body_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(body_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, body_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(continuing_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, continuing_, program_id);
 }
 
 LoopStatement::LoopStatement(LoopStatement&&) = default;

src/ast/matrix.cc

@@ -30,10 +30,10 @@ Matrix::Matrix(ProgramID program_id,
       subtype_(subtype),
       rows_(rows),
       columns_(columns) {
-  TINT_ASSERT(rows > 1);
+  TINT_ASSERT(AST, rows > 1);
-  TINT_ASSERT(rows < 5);
+  TINT_ASSERT(AST, rows < 5);
-  TINT_ASSERT(columns > 1);
+  TINT_ASSERT(AST, columns > 1);
-  TINT_ASSERT(columns < 5);
+  TINT_ASSERT(AST, columns < 5);
 }
 
 Matrix::Matrix(Matrix&&) = default;

src/ast/member_accessor_expression.cc

@@ -26,10 +26,10 @@ MemberAccessorExpression::MemberAccessorExpression(ProgramID program_id,
                                                    Expression* structure,
                                                    IdentifierExpression* member)
     : Base(program_id, source), struct_(structure), member_(member) {
-  TINT_ASSERT(struct_);
+  TINT_ASSERT(AST, struct_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(struct_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, struct_, program_id);
-  TINT_ASSERT(member_);
+  TINT_ASSERT(AST, member_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(member_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, member_, program_id);
 }
 
 MemberAccessorExpression::MemberAccessorExpression(MemberAccessorExpression&&) =

src/ast/module.cc

@@ -44,7 +44,7 @@ Module::Module(ProgramID program_id,
       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 @@ const ast::TypeDecl* Module::LookupType(Symbol name) const {
 }
 
 void Module::AddGlobalVariable(ast::Variable* var) {
-  TINT_ASSERT(var);
+  TINT_ASSERT(AST, var);
-  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);
   global_declarations_.push_back(var);
 }
 
 void Module::AddTypeDecl(ast::TypeDecl* type) {
-  TINT_ASSERT(type);
+  TINT_ASSERT(AST, type);
-  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);
   global_declarations_.push_back(type);
 }
 
 void Module::AddFunction(ast::Function* func) {
-  TINT_ASSERT(func);
+  TINT_ASSERT(AST, func);
-  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);
   global_declarations_.push_back(func);
 }
@@ -98,20 +98,22 @@ void Module::Copy(CloneContext* ctx, const Module* src) {
 
   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";
     }
   }
 }

src/ast/multisampled_texture.cc

@@ -26,7 +26,7 @@ MultisampledTexture::MultisampledTexture(ProgramID program_id,
                                          TextureDimension dim,
                                          Type* type)
     : Base(program_id, source, dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
 }
 
 MultisampledTexture::MultisampledTexture(MultisampledTexture&&) = default;

src/ast/return_statement.cc

@@ -28,7 +28,7 @@ ReturnStatement::ReturnStatement(ProgramID program_id,
                                  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;

src/ast/sampled_texture.cc

@@ -26,7 +26,7 @@ SampledTexture::SampledTexture(ProgramID program_id,
                                TextureDimension dim,
                                Type const* type)
     : Base(program_id, source, dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(AST, type_);
 }
 
 SampledTexture::SampledTexture(SampledTexture&&) = default;

src/ast/scalar_constructor_expression.cc

@@ -25,8 +25,8 @@ ScalarConstructorExpression::ScalarConstructorExpression(ProgramID program_id,
                                                          const Source& source,
                                                          Literal* literal)
     : Base(program_id, source), literal_(literal) {
-  TINT_ASSERT(literal);
+  TINT_ASSERT(AST, literal);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(literal, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, literal, program_id);
 }
 
 ScalarConstructorExpression::ScalarConstructorExpression(

src/ast/struct.cc

@@ -33,12 +33,12 @@ Struct::Struct(ProgramID program_id,
       members_(std::move(members)),
       decorations_(std::move(decorations)) {
   for (auto* mem : members_) {
-    TINT_ASSERT(mem);
+    TINT_ASSERT(AST, mem);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(mem, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, mem, program_id);
   }
   for (auto* deco : decorations_) {
-    TINT_ASSERT(deco);
+    TINT_ASSERT(AST, deco);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
 }
 

src/ast/struct_member.cc

@@ -30,12 +30,12 @@ StructMember::StructMember(ProgramID program_id,
       symbol_(sym),
       type_(type),
       decorations_(std::move(decorations)) {
-  TINT_ASSERT(type);
+  TINT_ASSERT(AST, type);
-  TINT_ASSERT(symbol_.IsValid());
+  TINT_ASSERT(AST, symbol_.IsValid());
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(symbol_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, symbol_, program_id);
   for (auto* deco : decorations_) {
-    TINT_ASSERT(deco);
+    TINT_ASSERT(AST, deco);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(deco, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, deco, program_id);
   }
 }
 

src/ast/switch_statement.cc

@@ -26,11 +26,11 @@ SwitchStatement::SwitchStatement(ProgramID program_id,
                                  Expression* condition,
                                  CaseStatementList body)
     : Base(program_id, source), condition_(condition), body_(body) {
-  TINT_ASSERT(condition_);
+  TINT_ASSERT(AST, condition_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(condition_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, condition_, program_id);
   for (auto* stmt : body_) {
-    TINT_ASSERT(stmt);
+    TINT_ASSERT(AST, stmt);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(stmt, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, stmt, program_id);
   }
 }
 

src/ast/type_constructor_expression.cc

@@ -26,10 +26,10 @@ TypeConstructorExpression::TypeConstructorExpression(ProgramID program_id,
                                                      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(AST, val);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(val, program_id);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, val, program_id);
   }
 }
 

src/ast/type_decl.cc

@@ -23,7 +23,7 @@ namespace ast {
 
 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;

src/ast/unary_op_expression.cc

@@ -26,8 +26,8 @@ UnaryOpExpression::UnaryOpExpression(ProgramID program_id,
                                      UnaryOp op,
                                      Expression* expr)
     : Base(program_id, source), op_(op), expr_(expr) {
-  TINT_ASSERT(expr_);
+  TINT_ASSERT(AST, expr_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(expr_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, expr_, program_id);
 }
 
 UnaryOpExpression::UnaryOpExpression(UnaryOpExpression&&) = default;

src/ast/variable.cc

@@ -40,9 +40,9 @@ Variable::Variable(ProgramID program_id,
       decorations_(std::move(decorations)),
       declared_storage_class_(declared_storage_class),
       declared_access_(declared_access) {
-  TINT_ASSERT(symbol_.IsValid());
+  TINT_ASSERT(AST, symbol_.IsValid());
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(symbol_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, symbol_, program_id);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(constructor, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, constructor, program_id);
 }
 
 Variable::Variable(Variable&&) = default;

src/ast/variable_decl_statement.cc

@@ -25,8 +25,8 @@ VariableDeclStatement::VariableDeclStatement(ProgramID program_id,
                                              const Source& source,
                                              Variable* variable)
     : Base(program_id, source), variable_(variable) {
-  TINT_ASSERT(variable_);
+  TINT_ASSERT(AST, variable_);
-  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(variable_, program_id);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, variable_, program_id);
 }
 
 VariableDeclStatement::VariableDeclStatement(VariableDeclStatement&&) = default;

src/ast/vector.cc

@@ -26,8 +26,8 @@ Vector::Vector(ProgramID program_id,
                Type const* subtype,
                uint32_t size)
     : Base(program_id, source), subtype_(subtype), size_(size) {
-  TINT_ASSERT(size_ > 1);
+  TINT_ASSERT(AST, size_ > 1);
-  TINT_ASSERT(size_ < 5);
+  TINT_ASSERT(AST, size_ < 5);
 }
 
 Vector::Vector(Vector&&) = default;

src/clone_context.h

@@ -105,7 +105,7 @@ class CloneContext {
     }
 
     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 @@ class CloneContext {
 
     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 @@ class CloneContext {
     }
 
     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 @@ class CloneContext {
     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 @@ class CloneContext {
   /// @returns this CloneContext so calls can be chained
   CloneContext& ReplaceAll(const SymbolTransform& replacer) {
     if (symbol_transform_) {
-      TINT_ICE(Diagnostics()) << "ReplaceAll(const SymbolTransform&) called "
+      TINT_ICE(Clone, Diagnostics())
-                                 "multiple times on the same CloneContext";
+          << "ReplaceAll(const SymbolTransform&) called "
+             "multiple times on the same CloneContext";
       return *this;
     }
     symbol_transform_ = replacer;
@@ -369,8 +370,8 @@ class CloneContext {
   /// @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(Clone, src, what);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, with);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(Clone, dst, with);
     replacements_[what] = with;
     return *this;
   }
@@ -382,9 +383,9 @@ class CloneContext {
   /// @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 @@ class CloneContext {
   /// @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 @@ class CloneContext {
   /// @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 @@ class CloneContext {
   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(Clone, src, before);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    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 @@ class CloneContext {
   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(Clone, src, after);
-    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(dst, object);
+    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 @@ class CloneContext {
     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;

src/clone_context_test.cc

@@ -783,7 +783,7 @@ TEST_F(CloneContextTest, ProgramIDs_Clone_ObjectNotOwnedBySrc) {
         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 @@ TEST_F(CloneContextTest, ProgramIDs_Clone_ObjectNotOwnedByDst) {
         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

src/debug.cc

@@ -27,12 +27,13 @@ void SetInternalCompilerErrorReporter(InternalCompilerErrorReporter* reporter) {
 
 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_);

src/debug.h

@@ -44,8 +44,12 @@ class InternalCompilerError {
   /// 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 @@ class InternalCompilerError {
  private:
   char const* const file_;
   size_t const line_;
+  diag::System system_;
   diag::List& diagnostics_;
   std::stringstream msg_;
 };
@@ -76,8 +81,9 @@ class InternalCompilerError {
 /// 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) \
+#define TINT_ICE(system, diagnostics)             \
-  tint::InternalCompilerError(__FILE__, __LINE__, 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 @@ class InternalCompilerError {
 /// 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) \
+#define TINT_UNREACHABLE(system, diagnostics) \
-  TINT_ICE(diagnostics) << "TINT_UNREACHABLE "
+  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 @@ class InternalCompilerError {
 /// 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) \
+#define TINT_UNIMPLEMENTED(system, diagnostics) \
-  TINT_ICE(diagnostics) << "TINT_UNIMPLEMENTED "
+  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 @@ class InternalCompilerError {
 /// 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)                                      \
+#define TINT_ASSERT(system, condition)                                      \
-  do {                                                              \
+  do {                                                                      \
-    if (!(condition)) {                                             \
+    if (!(condition)) {                                                     \
-      tint::diag::List diagnostics;                                 \
+      tint::diag::List diagnostics;                                         \
-      TINT_ICE(diagnostics) << "TINT_ASSERT(" << #condition << ")"; \
+      TINT_ICE(system, diagnostics) << "TINT_ASSERT(" << #condition << ")"; \
-    }                                                               \
+    }                                                                       \
   } while (false)
 
 #endif  // SRC_DEBUG_H_

src/debug_test.cc

@@ -23,17 +23,18 @@ TEST(DebugTest, Unreachable) {
   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

src/diagnostic/diagnostic.h

@@ -32,6 +32,24 @@ inline bool operator>=(Severity a, Severity b) {
   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 @@ class Diagnostic {
   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 @@ class List {
   }
 
   /// 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) {
+  void add_note(System system,
-    diag::Diagnostic error{};
+                const std::string& note_msg,
-    error.severity = diag::Severity::Note;
+                const Source& source) {
-    error.source = source;
+    diag::Diagnostic note{};
-    error.message = note_msg;
+    note.severity = diag::Severity::Note;
-    add(std::move(error));
+    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) {
+  void add_warning(System system,
-    diag::Diagnostic error{};
+                   const std::string& warning_msg,
-    error.severity = diag::Severity::Warning;
+                   const Source& source) {
-    error.source = source;
+    diag::Diagnostic warning{};
-    error.message = warning_msg;
+    warning.severity = diag::Severity::Warning;
-    add(std::move(error));
+    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 @@ class List {
 
   /// 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));

src/diagnostic/formatter_test.cc

@@ -33,18 +33,19 @@ class DiagFormatterTest : public testing::Test {
   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, SimpleNewlineAtEnd) {
 
 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 @@ the  snake  says  quack
 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 @@ the    snake    says    quack
 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 @@ the  snail  says  ???
 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
 

src/inspector/inspector.cc

@@ -47,7 +47,7 @@ namespace {
 
 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 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
     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 @@ void Inspector::AddEntryPointInOutVariables(
   }
 
   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();
 

src/intrinsic_table.cc

@@ -120,11 +120,11 @@ class ClosedState {
   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 @@ class ClosedState {
   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 @@ const sem::Intrinsic* Impl::Lookup(sem::IntrinsicType intrinsic_type,
       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 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
     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;
     }

src/program.cc

@@ -63,7 +63,7 @@ Program::Program(ProgramBuilder&& builder) {
     // 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 @@ std::string Program::str(const ast::Node* node) const {
 }
 
 void Program::AssertNotMoved() const {
-  TINT_ASSERT(!moved_);
+  TINT_ASSERT(Program, !moved_);
 }
 
 }  // namespace tint

src/program_builder.cc

@@ -83,7 +83,7 @@ void ProgramBuilder::MarkAsMoved() {
 
 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";
   }
 }

src/program_id.h

@@ -89,6 +89,7 @@ template <typename A, typename B>
 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 @@ void AssertProgramIDsEqual(A&& a,
     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
+/// TINT_ASSERT_PROGRAM_IDS_EQUAL(SYSTEM, A, B) is a macro that asserts that the
-/// identifiers for A and B are equal.
+/// program identifiers for A and B are equal.
 ///
-/// TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(A, B) is a macro that asserts that
+/// TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(SYSTEM, A, B) is a macro that asserts
-/// the program identifiers for A and B are equal, if both A and B have valid
+/// that the program identifiers for A and B are equal, if both A and B have
-/// program identifiers.
+/// valid program identifiers.
 #if TINT_CHECK_FOR_CROSS_PROGRAM_LEAKS
-#define TINT_ASSERT_PROGRAM_IDS_EQUAL(a, b)                                   \
+#define TINT_ASSERT_PROGRAM_IDS_EQUAL(system, a, b)                      \
-  detail::AssertProgramIDsEqual(                                              \
+  detail::AssertProgramIDsEqual(a, b, false, tint::diag::System::system, \
-      a, b, false, "TINT_ASSERT_PROGRAM_IDS_EQUAL(" #a ", " #b ")", __FILE__, \
+                                "TINT_ASSERT_PROGRAM_IDS_EQUAL(" #system \
-      __LINE__)
+                                "," #a ", " #b ")",                      \
-#define TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(a, b)                        \
+                                __FILE__, __LINE__)
-  detail::AssertProgramIDsEqual(                                            \
+#define TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(system, a, b)                 \
-      a, b, true, "TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(" #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) \

src/program_test.cc

@@ -87,7 +87,7 @@ TEST_F(ProgramTest, Assert_Null_Function) {
 }
 
 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());

src/reader/spirv/function.cc

@@ -779,7 +779,7 @@ FunctionEmitter::StatementBlock::StatementBlock(StatementBlock&& other) =
 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::Finalize(ProgramBuilder* pb) {
 }
 
 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::PushNewStatementBlock(const Construct* construct,
 
 void FunctionEmitter::PushGuard(const std::string& guard_name,
                                 uint32_t end_id) {
-  TINT_ASSERT(!statements_stack_.empty());
+  TINT_ASSERT(Reader, !statements_stack_.empty());
-  TINT_ASSERT(!guard_name.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::PushGuard(const std::string& guard_name,
 }
 
 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 @@ void FunctionEmitter::PushTrueGuard(uint32_t end_id) {
 }
 
 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::AddStatement(ast::Statement* statement) {
 }
 
 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 FunctionEmitter::Emit() {
 
   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 @@ bool FunctionEmitter::Emit() {
 }
 
 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 @@ bool FunctionEmitter::EmitEntryPointAsWrapper() {
   // 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(Reader, var != nullptr);
-    TINT_ASSERT(var->opcode() == SpvOpVariable);
+    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 @@ bool FunctionEmitter::EmitEntryPointAsWrapper() {
             create<ast::BuiltinDecoration>(source, ast::Builtin::kPosition));
       } else {
         const auto* var = def_use_mgr_->GetDef(var_id);
-        TINT_ASSERT(var != nullptr);
+        TINT_ASSERT(Reader, var != nullptr);
-        TINT_ASSERT(var->opcode() == SpvOpVariable);
+        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 @@ bool FunctionEmitter::LabelControlFlowConstructs() {
   //      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 @@ bool FunctionEmitter::LabelControlFlowConstructs() {
     // 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(Reader, parent);
-      TINT_ASSERT(parent->kind == Construct::kContinue);
+      TINT_ASSERT(Reader, parent->kind == Construct::kContinue);
       scope_end_pos = parent->end_pos;
       parent = parent->parent;
     }
@@ -1543,9 +1544,9 @@ bool FunctionEmitter::LabelControlFlowConstructs() {
 
   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 @@ bool FunctionEmitter::LabelControlFlowConstructs() {
       }
     }
 
-    TINT_ASSERT(top);
+    TINT_ASSERT(Reader, top);
     block_info->construct = top;
   }
 
@@ -1828,9 +1829,9 @@ bool FunctionEmitter::ClassifyCFGEdges() {
   //    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 @@ bool FunctionEmitter::ClassifyCFGEdges() {
     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 @@ bool FunctionEmitter::ClassifyCFGEdges() {
                         << " (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 @@ bool FunctionEmitter::FindIfSelectionInternalHeaders() {
       // 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 @@ bool FunctionEmitter::EmitFunctionBodyStatements() {
 
   // 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(Reader, function_construct != nullptr);
-  TINT_ASSERT(function_construct->kind == Construct::kFunction);
+  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::EmitBasicBlock(const BlockInfo& block_info) {
 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(Reader, construct->kind == Construct::kIfSelection);
-  TINT_ASSERT(construct->begin_id == block_info.id);
+  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::EmitIfStart(const BlockInfo& block_info) {
 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(Reader, construct->kind == Construct::kSwitchSelection);
-  TINT_ASSERT(construct->begin_id == block_info.id);
+  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 @@ bool FunctionEmitter::EmitSwitchStart(const BlockInfo& block_info) {
       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 @@ ast::Statement* FunctionEmitter::MakeBranchDetailed(
       // 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 @@ bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info,
   // 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 @@ bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info,
   // 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 @@ bool FunctionEmitter::EmitStatement(const spvtools::opt::Instruction& inst) {
           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 @@ TypedExpression FunctionEmitter::MakeAccessChain(
     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 @@ TypedExpression FunctionEmitter::MakeVectorShuffle(
           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 @@ const Construct* FunctionEmitter::GetEnclosingScope(uint32_t first_pos,
                                                     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 @@ const Construct* FunctionEmitter::GetEnclosingScope(uint32_t first_pos,
         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 @@ ast::Expression* FunctionEmitter::ConvertTexelForStorage(
 
   // The texel type is always a 4-element vector.
   const uint32_t dest_count = 4u;
-  TINT_ASSERT(dest_type->Is<Vector>() &&
+  TINT_ASSERT(Reader, dest_type->Is<Vector>() &&
-              dest_type->As<Vector>()->size == dest_count);
+                          dest_type->As<Vector>()->size == dest_count);
-  TINT_ASSERT(dest_type->IsFloatVector() ||
+  TINT_ASSERT(Reader, dest_type->IsFloatVector() ||
-              dest_type->IsUnsignedIntegerVector() ||
+                          dest_type->IsUnsignedIntegerVector() ||
-              dest_type->IsSignedIntegerVector());
+                          dest_type->IsSignedIntegerVector());
 
   if (src_type == dest_type) {
     return texel.expr;
@@ -5294,7 +5296,7 @@ ast::Expression* FunctionEmitter::ConvertTexelForStorage(
   }
 
   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;

src/reader/spirv/function.h

@@ -1017,7 +1017,7 @@ class FunctionEmitter {
   /// @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)...);
   }

src/reader/spirv/namer.cc

@@ -123,7 +123,7 @@ std::string Namer::FindUnusedDerivedName(const std::string& base_name) const {
 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;
 }
 

src/reader/spirv/parser.cc

@@ -29,7 +29,7 @@ Program Parse(const std::vector<uint32_t>& input) {
   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));
   }
 

src/reader/spirv/parser_impl.cc

@@ -833,8 +833,8 @@ bool ParserImpl::RegisterEntryPoints() {
       // 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(Reader, !inner_implementation_name.empty());
-    TINT_ASSERT(ep_name != inner_implementation_name);
+    TINT_ASSERT(Reader, ep_name != inner_implementation_name);
 
     tint::UniqueVector<uint32_t> inputs;
     tint::UniqueVector<uint32_t> outputs;

src/reader/spirv/parser_impl_barrier_test.cc

@@ -40,7 +40,7 @@ Program ParseAndBuild(std::string spirv) {
   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();

src/reader/wgsl/parser_impl.cc

@@ -269,14 +269,15 @@ ParserImpl::Failure::Errored ParserImpl::add_error(const Token& t,
 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 @@ T ParserImpl::sync(Token::Type tok, F&& body) {
   --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();
 

src/reader/wgsl/parser_impl.h

@@ -112,7 +112,7 @@ class ParserImpl {
     /// 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 @@ class ParserImpl {
     /// 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);
     }
 

src/resolver/resolver.h

@@ -409,6 +409,21 @@ class Resolver {
   /// @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,

src/sem/array.cc

@@ -35,7 +35,7 @@ Array::Array(const Type* element,
       size_(size),
       stride_(stride),
       implicit_stride_(implicit_stride) {
-  TINT_ASSERT(element_);
+  TINT_ASSERT(Semantic, element_);
 }
 
 std::string Array::type_name() const {

src/sem/atomic_type.cc

@@ -23,7 +23,7 @@ namespace tint {
 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 {

src/sem/call_target.cc

@@ -23,7 +23,7 @@ namespace sem {
 
 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;

src/sem/depth_texture_type.cc

@@ -32,7 +32,7 @@ bool IsValidDepthDimension(ast::TextureDimension dim) {
 }  // namespace
 
 DepthTexture::DepthTexture(ast::TextureDimension dim) : Base(dim) {
-  TINT_ASSERT(IsValidDepthDimension(dim));
+  TINT_ASSERT(Semantic, IsValidDepthDimension(dim));
 }
 
 DepthTexture::DepthTexture(DepthTexture&&) = default;

src/sem/expression.cc

@@ -23,7 +23,7 @@ Expression::Expression(const ast::Expression* declaration,
                        const sem::Type* type,
                        Statement* statement)
     : declaration_(declaration), type_(type), statement_(statement) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 }  // namespace sem

src/sem/info.h

@@ -70,7 +70,7 @@ class Info {
   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);
   }
 

src/sem/matrix_type.cc

@@ -27,10 +27,10 @@ Matrix::Matrix(Vector* column_type, uint32_t columns)
       column_type_(column_type),
       rows_(column_type->size()),
       columns_(columns) {
-  TINT_ASSERT(rows_ > 1);
+  TINT_ASSERT(AST, rows_ > 1);
-  TINT_ASSERT(rows_ < 5);
+  TINT_ASSERT(AST, rows_ < 5);
-  TINT_ASSERT(columns_ > 1);
+  TINT_ASSERT(AST, columns_ > 1);
-  TINT_ASSERT(columns_ < 5);
+  TINT_ASSERT(AST, columns_ < 5);
 }
 
 Matrix::Matrix(Matrix&&) = default;

src/sem/multisampled_texture_type.cc

@@ -24,7 +24,7 @@ namespace sem {
 MultisampledTexture::MultisampledTexture(ast::TextureDimension dim,
                                          const Type* type)
     : Base(dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 MultisampledTexture::MultisampledTexture(MultisampledTexture&&) = default;

src/sem/pointer_type.cc

@@ -26,8 +26,8 @@ Pointer::Pointer(const Type* subtype,
                  ast::StorageClass storage_class,
                  ast::Access access)
     : subtype_(subtype), storage_class_(storage_class), access_(access) {
-  TINT_ASSERT(!subtype->Is<Reference>());
+  TINT_ASSERT(Semantic, !subtype->Is<Reference>());
-  TINT_ASSERT(access != ast::Access::kUndefined);
+  TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
 std::string Pointer::type_name() const {

src/sem/reference_type.cc

@@ -25,8 +25,8 @@ Reference::Reference(const Type* subtype,
                      ast::StorageClass storage_class,
                      ast::Access access)
     : subtype_(subtype), storage_class_(storage_class), access_(access) {
-  TINT_ASSERT(!subtype->Is<Reference>());
+  TINT_ASSERT(Semantic, !subtype->Is<Reference>());
-  TINT_ASSERT(access != ast::Access::kUndefined);
+  TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
 std::string Reference::type_name() const {

src/sem/sampled_texture_type.cc

@@ -23,7 +23,7 @@ namespace sem {
 
 SampledTexture::SampledTexture(ast::TextureDimension dim, const Type* type)
     : Base(dim), type_(type) {
-  TINT_ASSERT(type_);
+  TINT_ASSERT(Semantic, type_);
 }
 
 SampledTexture::SampledTexture(SampledTexture&&) = default;

src/sem/statement.cc

@@ -47,7 +47,7 @@ Statement::Statement(const ast::Statement* declaration, const Statement* parent)
             }
           }
         }
-        TINT_ASSERT(statement_is_continuing_for_loop);
+        TINT_ASSERT(Semantic, statement_is_continuing_for_loop);
       }
     }
   }

src/sem/vector_type.cc

@@ -23,8 +23,8 @@ namespace sem {
 
 Vector::Vector(Type const* subtype, uint32_t size)
     : subtype_(subtype), size_(size) {
-  TINT_ASSERT(size_ > 1);
+  TINT_ASSERT(Semantic, size_ > 1);
-  TINT_ASSERT(size_ < 5);
+  TINT_ASSERT(Semantic, size_ < 5);
 }
 
 Vector::Vector(Vector&&) = default;

src/symbol.cc

@@ -32,7 +32,8 @@ Symbol& Symbol::operator=(const Symbol& o) = default;
 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_;
 }
 

src/symbol_table.cc

@@ -54,7 +54,7 @@ Symbol SymbolTable::Get(const std::string& name) const {
 }
 
 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();

src/transform/array_length_from_uniform.cc

@@ -38,6 +38,7 @@ Output ArrayLengthFromUniform::Run(const Program* in, const DataMap& data) {
   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 @@ Output ArrayLengthFromUniform::Run(const Program* in, const DataMap& data) {
     // 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 @@ Output ArrayLengthFromUniform::Run(const Program* in, const DataMap& data) {
     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 @@ Output ArrayLengthFromUniform::Run(const Program* in, const DataMap& data) {
     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.group) + "," +
-          std::to_string(binding.binding) + ")");
+              std::to_string(binding.binding) + ")");
       continue;
     }
 

src/transform/binding_remapper.cc

@@ -45,6 +45,7 @@ Output BindingRemapper::Run(const Program* in, const DataMap& datamap) {
   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)));
   }

src/transform/bound_array_accessors.cc

@@ -74,7 +74,8 @@ ast::ArrayAccessorExpression* BoundArrayAccessors::Transform(
       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 @@ ast::ArrayAccessorExpression* BoundArrayAccessors::Transform(
     } 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;
     }

src/transform/calculate_array_length.cc

@@ -134,7 +134,7 @@ Output CalculateArrayLength::Run(const Program* in, const DataMap&) {
           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 @@ Output CalculateArrayLength::Run(const Program* in, const DataMap&) {
 
           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 @@ Output CalculateArrayLength::Run(const Program* in, const DataMap&) {
           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;

src/transform/canonicalize_entry_point_io.cc

@@ -69,6 +69,7 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap& data) {
   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 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap& data) {
           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 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap& data) {
         // 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(

src/transform/decompose_memory_access.cc

@@ -352,7 +352,7 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicAtomicFor(ProgramBuilder* builder,
       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 @@ struct DecomposeMemoryAccess::State {
   /// @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 @@ struct DecomposeMemoryAccess::State {
 
       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();
       }

src/transform/external_texture_transform.cc

@@ -58,7 +58,7 @@ Output ExternalTextureTransform::Run(const Program* in, const DataMap&) {
                     ->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 @@ Output ExternalTextureTransform::Run(const Program* in, const DataMap&) {
               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";

src/transform/fold_constants.cc

@@ -67,7 +67,7 @@ struct Value {
   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 @@ struct Value {
         return func(elems[index].bool_);
       }
     }
-    TINT_ASSERT(false && "Unreachable");
+    TINT_ASSERT(Transform, false && "Unreachable");
     return func(~0);
   }
 };
@@ -105,7 +105,7 @@ Value::Type AstToValueType(ast::Type* t) {
   } 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 @@ Value Fold(const ast::ScalarConstructorExpression* scalar_ctor) {
   if (auto* lit = literal->As<ast::BoolLiteral>()) {
     return {lit->IsTrue()};
   }
-  TINT_ASSERT(false && "Unreachable");
+  TINT_ASSERT(Transform, false && "Unreachable");
   return {};
 }
 

src/transform/promote_initializers_to_const_var.cc

@@ -52,7 +52,7 @@ Output PromoteInitializersToConstVar::Run(const Program* in, const DataMap&) {
     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;
       }

src/transform/renamer.cc

@@ -857,7 +857,7 @@ Output Renamer::Run(const Program* in, const DataMap&) {
     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 @@ Output Renamer::Run(const Program* in, const DataMap&) {
     } 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>()) {

src/transform/single_entry_point.cc

@@ -35,7 +35,8 @@ Output SingleEntryPoint::Run(const Program* in, const DataMap& data) {
 
   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 @@ Output SingleEntryPoint::Run(const Program* in, const DataMap& data) {
     }
   }
   if (entry_point == nullptr) {
-    out.Diagnostics().add_error("entry point '" + cfg->entry_point_name +
+    out.Diagnostics().add_error(
-                                "' not found");
+        diag::System::Transform,
+        "entry point '" + cfg->entry_point_name + "' not found");
     return Output(Program(std::move(out)));
   }
 
@@ -81,7 +83,7 @@ Output SingleEntryPoint::Run(const Program* in, const DataMap& data) {
         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)));
     }

src/transform/transform.cc

@@ -134,7 +134,7 @@ ast::Type* Transform::CreateASTTypeFor(CloneContext* ctx, const sem::Type* ty) {
   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;
 }

src/transform/vertex_pulling.cc

@@ -304,7 +304,8 @@ struct State {
       }
       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 @@ struct State {
         }
         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 @@ Output VertexPulling::Run(const Program* in, const DataMap& data) {
   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)));
   }
 

src/transform/zero_init_workgroup_memory.cc

@@ -78,7 +78,7 @@ struct ZeroInitWorkgroupMemory::State {
       return;
     }
 
-    TINT_UNREACHABLE(ctx.dst->Diagnostics())
+    TINT_UNREACHABLE(Transform, ctx.dst->Diagnostics())
         << "could not zero workgroup type: " << ty->type_name();
   }
 

src/utils/io/tmpfile_posix.cc

@@ -30,8 +30,8 @@ std::string TmpFilePath(std::string ext) {
   // (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() <=
+  TINT_ASSERT(Utils, ext.length() <=
-              static_cast<size_t>(std::numeric_limits<int>::max()));
+                         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) {

src/writer/append_vector.cc

@@ -60,8 +60,9 @@ ast::TypeConstructorExpression* AppendVector(ProgramBuilder* b,
   } 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: "
+    TINT_UNREACHABLE(Writer, b->Diagnostics())
-                                       << packed_el_sem_ty->TypeInfo().name;
+        << "unsupported vector element type: "
+        << packed_el_sem_ty->TypeInfo().name;
   }
 
   auto* statement = vector_sem->Stmt();
@@ -86,7 +87,7 @@ ast::TypeConstructorExpression* AppendVector(ProgramBuilder* b,
         } 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;
         }

src/writer/float_to_string.cc

@@ -106,7 +106,7 @@ std::string FloatToBitPreservingString(float f) {
         }
       } 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

src/writer/hlsl/generator_impl.cc

@@ -152,7 +152,7 @@ bool GeneratorImpl::Generate(std::ostream& out) {
         }
       }
     } else {
-      TINT_ICE(diagnostics_)
+      TINT_ICE(Writer, diagnostics_)
           << "unhandled module-scope declaration: " << decl->TypeInfo().name;
       return false;
     }
@@ -188,7 +188,8 @@ bool GeneratorImpl::EmitBitcast(std::ostream& pre,
                                 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 @@ bool GeneratorImpl::EmitBinary(std::ostream& pre,
     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 @@ bool GeneratorImpl::EmitBinary(std::ostream& pre,
       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 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
         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 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
 
   auto* func = builder_.AST().Functions().Find(ident->symbol());
   if (func == nullptr) {
-    diagnostics_.add_error("Unable to find function: " +
+    diagnostics_.add_error(diag::System::Writer,
-                           builder_.Symbols().NameFor(ident->symbol()));
+                           "Unable to find function: " +
+                               builder_.Symbols().NameFor(ident->symbol()));
     return false;
   }
 
@@ -610,7 +613,7 @@ bool GeneratorImpl::EmitUniformBufferAccess(
         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 @@ bool GeneratorImpl::EmitUniformBufferAccess(
     default:
       break;
   }
-  TINT_UNREACHABLE(diagnostics_)
+  TINT_UNREACHABLE(Writer, diagnostics_)
       << "unsupported DecomposeMemoryAccess::Intrinsic::Op: "
       << static_cast<int>(intrinsic->op);
   return false;
@@ -681,7 +684,7 @@ bool GeneratorImpl::EmitStorageBufferAccess(
         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 @@ bool GeneratorImpl::EmitStorageBufferAccess(
         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 @@ bool GeneratorImpl::EmitStorageBufferAccess(
       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 @@ bool GeneratorImpl::EmitStorageAtomicCall(
       break;
 
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
           << static_cast<int>(op);
       return false;
@@ -1077,7 +1080,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& pre,
       break;
 
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "unsupported atomic intrinsic: " << intrinsic->Type();
       return false;
   }
@@ -1266,6 +1269,7 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
       break;
     default:
       diagnostics_.add_error(
+          diag::System::Writer,
           "Internal error: unhandled data packing intrinsic");
       return false;
   }
@@ -1337,6 +1341,7 @@ bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& pre,
       break;
     default:
       diagnostics_.add_error(
+          diag::System::Writer,
           "Internal error: unhandled data packing intrinsic");
       return false;
   }
@@ -1354,7 +1359,7 @@ bool GeneratorImpl::EmitBarrierCall(std::ostream&,
   } 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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
 
   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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
         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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
         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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
         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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
         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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
           }
           break;
         default:
-          TINT_ICE(diagnostics_) << "unexpected intrinsic";
+          TINT_ICE(Writer, diagnostics_) << "unexpected intrinsic";
           return false;
       }
 
@@ -1496,7 +1502,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
       }
 
       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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
       } 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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
       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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
 
   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 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
       }
     }
     if (wgsl_ret_width > hlsl_ret_width) {
-      TINT_ICE(diagnostics_) << "WGSL return width (" << wgsl_ret_width
+      TINT_ICE(Writer, diagnostics_)
-                             << ") is wider than HLSL return width ("
+          << "WGSL return width (" << wgsl_ret_width
-                             << hlsl_ret_width << ") for " << intrinsic->Type();
+          << ") is wider than HLSL return width (" << hlsl_ret_width << ") for "
+          << intrinsic->Type();
       return false;
     }
   }
@@ -1766,8 +1775,9 @@ std::string GeneratorImpl::generate_builtin_name(
     case sem::IntrinsicType::kSmoothStep:
       return "smoothstep";
     default:
-      diagnostics_.add_error("Unknown builtin method: " +
+      diagnostics_.add_error(
-                             std::string(intrinsic->str()));
+          diag::System::Writer,
+          "Unknown builtin method: " + std::string(intrinsic->str()));
   }
 
   return "";
@@ -1936,7 +1946,8 @@ bool GeneratorImpl::EmitExpression(std::ostream& pre,
     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 @@ bool GeneratorImpl::EmitGlobalVariable(std::ostream& out,
       break;
   }
 
-  TINT_ICE(diagnostics_) << "unhandled storage class " << sem->StorageClass();
+  TINT_ICE(Writer, diagnostics_)
+      << "unhandled storage class " << sem->StorageClass();
   return false;
 }
 
@@ -2099,7 +2111,7 @@ bool GeneratorImpl::EmitUniformVariable(std::ostream& out,
   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 @@ bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out,
       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 @@ bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out,
     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 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, ast::Literal* lit) {
   } 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 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
     }
     out << "}";
   } else {
-    diagnostics_.add_error("Invalid type for zero emission: " +
+    diagnostics_.add_error(
-                           type->type_name());
+        diag::System::Writer,
+        "Invalid type for zero emission: " + type->type_name());
     return false;
   }
   return true;
@@ -2556,7 +2570,8 @@ bool GeneratorImpl::EmitStatement(std::ostream& out, ast::Statement* stmt) {
     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 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       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 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     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 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     // 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 @@ bool GeneratorImpl::EmitType(std::ostream& out,
         out << "CubeArray";
         break;
       default:
-        TINT_UNREACHABLE(diagnostics_)
+        TINT_UNREACHABLE(Writer, diagnostics_)
             << "unexpected TextureDimension " << tex->dim();
         return false;
     }
@@ -2712,8 +2727,9 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     if (storage) {
       auto* component = image_format_to_rwtexture_type(storage->image_format());
       if (component == nullptr) {
-        TINT_ICE(diagnostics_) << "Unsupported StorageTexture ImageFormat: "
+        TINT_ICE(Writer, diagnostics_)
-                               << static_cast<int>(storage->image_format());
+            << "Unsupported StorageTexture ImageFormat: "
+            << static_cast<int>(storage->image_format());
         return false;
       }
       out << "<" << component << ">";
@@ -2727,7 +2743,8 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       } 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 @@ bool GeneratorImpl::EmitType(std::ostream& out,
   } 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 @@ bool GeneratorImpl::EmitStructType(std::ostream& out, const sem::Struct* str) {
       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 @@ bool GeneratorImpl::EmitStructType(std::ostream& out, const sem::Struct* str) {
                        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 @@ bool GeneratorImpl::EmitVariable(std::ostream& out, ast::Variable* var) {
 
   // 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 @@ bool GeneratorImpl::EmitProgramConstVariable(std::ostream& out,
 
   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;
   }
 

src/writer/msl/generator_impl.cc

@@ -101,7 +101,7 @@ bool GeneratorImpl::Generate() {
         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 @@ bool GeneratorImpl::EmitTypeDecl(const sem::Type* ty) {
       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 @@ bool GeneratorImpl::EmitBinary(ast::BinaryExpression* expr) {
       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 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
 
   auto* func = program_->AST().Functions().Find(ident->symbol());
   if (func == nullptr) {
-    diagnostics_.add_error("Unable to find function: " +
+    diagnostics_.add_error(diag::System::Writer,
-                           program_->Symbols().NameFor(ident->symbol()));
+                           "Unable to find function: " +
+                               program_->Symbols().NameFor(ident->symbol()));
     return false;
   }
 
@@ -425,7 +428,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
 
   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 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
       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 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
       out_ << ".write(";
       break;
     default:
-      TINT_UNREACHABLE(diagnostics_)
+      TINT_UNREACHABLE(Writer, diagnostics_)
           << "Unhandled texture intrinsic '" << intrinsic->str() << "'";
       return false;
   }
@@ -595,7 +599,8 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
             out_ << "uint3(";
             break;
           default:
-            TINT_ICE(diagnostics_) << "unhandled texture dimensionality";
+            TINT_ICE(Writer, diagnostics_)
+                << "unhandled texture dimensionality";
             break;
         }
       }
@@ -653,7 +658,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
       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 @@ std::string GeneratorImpl::generate_builtin_name(
       out += "unpack_unorm2x16_to_float";
       break;
     default:
-      diagnostics_.add_error("Unknown import method: " +
+      diagnostics_.add_error(
-                             std::string(intrinsic->str()));
+          diag::System::Writer,
+          "Unknown import method: " + std::string(intrinsic->str()));
       return "";
   }
   return out;
@@ -949,8 +955,9 @@ bool GeneratorImpl::EmitZeroValue(const sem::Type* type) {
   } else if (type->As<sem::Struct>()) {
     out_ << "{}";
   } else {
-    diagnostics_.add_error("Invalid type for zero emission: " +
+    diagnostics_.add_error(
-                           type->type_name());
+        diag::System::Writer,
+        "Invalid type for zero emission: " + type->type_name());
     return false;
   }
   return true;
@@ -971,7 +978,7 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
   } 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 @@ bool GeneratorImpl::EmitExpression(ast::Expression* expr) {
     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 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
       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 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
       } 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 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
 
         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 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     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 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
     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 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
         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 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
         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 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
       } 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 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
       }
       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 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
   } 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 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
     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 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
       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 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
                        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 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
       // 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 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var) {
 
   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 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var) {
       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 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
 
   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 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
 
   if (auto* arr = ty->As<sem::Array>()) {
     if (!arr->IsStrideImplicit()) {
-      TINT_ICE(diagnostics_) << "arrays with explicit strides should have "
+      TINT_ICE(Writer, diagnostics_)
-                                "removed with the PadArrayElements transform";
+          << "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 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
     return SizeAndAlign{str->Size(), str->Align()};
   }
 
-  TINT_UNREACHABLE(diagnostics_) << "Unhandled type " << ty->TypeInfo().name;
+  TINT_UNREACHABLE(Writer, diagnostics_)
+      << "Unhandled type " << ty->TypeInfo().name;
   return {};
 }
 

src/writer/spirv/builder.cc

@@ -453,7 +453,7 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
       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 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
           // 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 @@ bool Builder::GenerateArrayAccessor(ast::ArrayAccessorExpression* expr,
   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 @@ bool Builder::GenerateArrayAccessor(ast::ArrayAccessorExpression* expr,
     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 @@ bool Builder::GenerateMemberAccessor(ast::MemberAccessorExpression* expr,
     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 @@ bool Builder::GenerateMemberAccessor(ast::MemberAccessorExpression* expr,
 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 @@ uint32_t Builder::GenerateCastOrCopyOrPassthrough(const sem::Type* to_type,
   // 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 @@ uint32_t Builder::GenerateCastOrCopyOrPassthrough(const sem::Type* to_type,
     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 @@ bool Builder::GenerateTextureIntrinsic(ast::CallExpression* call,
   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 @@ bool Builder::GenerateTextureIntrinsic(ast::CallExpression* call,
 
   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 @@ bool Builder::GenerateTextureIntrinsic(ast::CallExpression* call,
       break;
     }
     default:
-      TINT_UNREACHABLE(builder_.Diagnostics());
+      TINT_UNREACHABLE(Writer, builder_.Diagnostics());
       return false;
   }
 
@@ -2942,7 +2944,7 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
           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 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
         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 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
                                 });
     }
     default:
-      TINT_UNREACHABLE(builder_.Diagnostics())
+      TINT_UNREACHABLE(Writer, builder_.Diagnostics())
           << "unhandled atomic intrinsic " << intrinsic->Type();
       return false;
   }
@@ -3374,7 +3376,7 @@ bool Builder::GenerateSwitchStatement(ast::SwitchStatement* stmt) {
     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 @@ SpvBuiltIn Builder::ConvertBuiltin(ast::Builtin builtin,
       } 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:

src/writer/spirv/builder.h

@@ -192,7 +192,7 @@ class Builder {
   /// @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);

src/writer/wgsl/generator_impl.cc

@@ -80,7 +80,7 @@ bool GeneratorImpl::Generate() {
         return false;
       }
     } else {
-      TINT_UNREACHABLE(diagnostics_);
+      TINT_UNREACHABLE(Writer, diagnostics_);
       return false;
     }
 
@@ -106,8 +106,9 @@ bool GeneratorImpl::EmitTypeDecl(const ast::TypeDecl* ty) {
       return false;
     }
   } else {
-    diagnostics_.add_error("unknown declared type: " +
+    diagnostics_.add_error(
-                           std::string(ty->TypeInfo().name));
+        diag::System::Writer,
+        "unknown declared type: " + std::string(ty->TypeInfo().name));
     return false;
   }
   return true;
@@ -139,7 +140,7 @@ bool GeneratorImpl::EmitExpression(ast::Expression* expr) {
     return EmitUnaryOp(u);
   }
 
-  diagnostics_.add_error("unknown expression type");
+  diagnostics_.add_error(diag::System::Writer, "unknown expression type");
   return false;
 }
 
@@ -273,7 +274,7 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
   } 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::EmitFunction(ast::Function* func) {
 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 @@ bool GeneratorImpl::EmitAccess(const ast::Access access) {
     default:
       break;
   }
-  diagnostics_.add_error("unknown access");
+  diagnostics_.add_error(diag::System::Writer, "unknown access");
   return false;
 }
 
@@ -435,7 +436,7 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
     } 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 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
         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 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
   } 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 @@ bool GeneratorImpl::EmitDecorations(const ast::DecorationList& decos) {
               return false;
             }
           } else {
-            TINT_ICE(diagnostics_) << "Unsupported workgroup_size expression";
+            TINT_ICE(Writer, diagnostics_)
+                << "Unsupported workgroup_size expression";
           }
         }
       }
@@ -676,7 +680,7 @@ bool GeneratorImpl::EmitDecorations(const ast::DecorationList& decos) {
     } 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 @@ bool GeneratorImpl::EmitBinary(ast::BinaryExpression* expr) {
       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 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
     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 @@ std::string GeneratorImpl::UniqueIdentifier(const std::string& suffix) {
       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>";
 }