diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index 7ea3010..24ef2fd 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc
@@ -99,6 +99,12 @@
 #include "src/tint/lang/glsl/validate/validate.h"
 #endif  // TINT_BUILD_GLSL_VALIDATOR
 
+#if TINT_BUILD_WGSL_READER
+#define WGSL_READER_ONLY(x) x
+#else
+#define WGSL_READER_ONLY(x)
+#endif
+
 #if TINT_BUILD_SPV_WRITER
 #define SPV_WRITER_ONLY(x) x
 #else
@@ -145,6 +151,7 @@
     kMsl,
     kHlsl,
     kGlsl,
+    kIr,
 };
 
 #if TINT_BUILD_HLSL_WRITER
@@ -296,6 +303,7 @@
     MSL_WRITER_ONLY(format_enum_names.Emplace(Format::kMsl, "msl"));
     HLSL_WRITER_ONLY(format_enum_names.Emplace(Format::kHlsl, "hlsl"));
     GLSL_WRITER_ONLY(format_enum_names.Emplace(Format::kGlsl, "glsl"));
+    WGSL_READER_ONLY(format_enum_names.Emplace(Format::kIr, "ir"));
 
     OptionSet options;
     auto& fmt = options.Add<EnumOption<Format>>("format",
@@ -1199,6 +1207,27 @@
 #endif  // TINT_BUILD_GLSL_WRITER
 }
 
+/// Generate IR code for a program.
+/// @param program the program to generate
+/// @param options the options that Tint was invoked with
+/// @returns true on success
+bool GenerateIr([[maybe_unused]] const tint::Program& program,
+                [[maybe_unused]] const Options& options) {
+#if !TINT_BUILD_WGSL_READER
+    std::cerr << "WGSL reader not enabled in tint build" << std::endl;
+    return false;
+#else
+    auto result = tint::wgsl::reader::ProgramToLoweredIR(program);
+    if (result != tint::Success) {
+        std::cerr << "Failed to build IR from program: " << result.Failure() << "\n";
+        return false;
+    }
+    options.printer->Print(tint::core::ir::Disassemble(result.Get()));
+    options.printer->Print(tint::StyledText{} << "\n");
+    return true;
+#endif
+}
+
 }  // namespace
 
 int main(int argc, const char** argv) {
@@ -1336,15 +1365,7 @@
 
 #if TINT_BUILD_WGSL_READER
     if (options.dump_ir) {
-        auto result = tint::wgsl::reader::ProgramToLoweredIR(info.program);
-        if (result != tint::Success) {
-            std::cerr << "Failed to build IR from program: " << result.Failure() << "\n";
-        } else {
-            auto mod = result.Move();
-            if (options.dump_ir) {
-                std::cout << tint::core::ir::Disassemble(mod) << "\n";
-            }
-        }
+        GenerateIr(info.program, options);
     }
 #endif  // TINT_BUILD_WGSL_READER
 
@@ -1455,6 +1476,9 @@
         case Format::kGlsl:
             success = GenerateGlsl(program, options);
             break;
+        case Format::kIr:
+            success = GenerateIr(program, options);
+            break;
         case Format::kNone:
             break;
         default:
diff --git a/src/tint/lang/core/ir/binary/roundtrip_fuzz.cc b/src/tint/lang/core/ir/binary/roundtrip_fuzz.cc
index 32886de..ccc7a9c 100644
--- a/src/tint/lang/core/ir/binary/roundtrip_fuzz.cc
+++ b/src/tint/lang/core/ir/binary/roundtrip_fuzz.cc
@@ -46,8 +46,8 @@
         return;
     }
 
-    auto in = Disassemble(module);
-    auto out = Disassemble(decoded.Get());
+    auto in = Disassemble(module).Plain();
+    auto out = Disassemble(decoded.Get()).Plain();
     if (in != out) {
         TINT_ICE() << "Roundtrip produced different disassembly\n"
                    << "-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-\n"
diff --git a/src/tint/lang/core/ir/binary/roundtrip_test.cc b/src/tint/lang/core/ir/binary/roundtrip_test.cc
index 19344f1..9d8b368 100644
--- a/src/tint/lang/core/ir/binary/roundtrip_test.cc
+++ b/src/tint/lang/core/ir/binary/roundtrip_test.cc
@@ -47,7 +47,7 @@
 class IRBinaryRoundtripTestBase : public IRTestParamHelper<T> {
   public:
     std::pair<std::string, std::string> Roundtrip() {
-        auto pre = Disassemble(this->mod);
+        auto pre = Disassemble(this->mod).Plain();
         auto encoded = Encode(this->mod);
         if (encoded != Success) {
             return {pre, encoded.Failure().reason.Str()};
@@ -56,7 +56,7 @@
         if (decoded != Success) {
             return {pre, decoded.Failure().reason.Str()};
         }
-        auto post = Disassemble(decoded.Get());
+        auto post = Disassemble(decoded.Get()).Plain();
         return {pre, post};
     }
 };
diff --git a/src/tint/lang/core/ir/disassembler.cc b/src/tint/lang/core/ir/disassembler.cc
index c31cad8..1664dce 100644
--- a/src/tint/lang/core/ir/disassembler.cc
+++ b/src/tint/lang/core/ir/disassembler.cc
@@ -66,15 +66,32 @@
 #include "src/tint/lang/core/type/struct.h"
 #include "src/tint/lang/core/type/type.h"
 #include "src/tint/utils/ice/ice.h"
+#include "src/tint/utils/macros/defer.h"
 #include "src/tint/utils/macros/scoped_assignment.h"
 #include "src/tint/utils/rtti/switch.h"
 #include "src/tint/utils/text/string.h"
+#include "src/tint/utils/text/styled_text.h"
+#include "src/tint/utils/text/text_style.h"
 
 using namespace tint::core::fluent_types;  // NOLINT
 
 namespace tint::core::ir {
 namespace {
 
+static constexpr auto StylePlain = style::Plain;
+static constexpr auto StyleAttribute = style::Attribute + style::NoQuote;
+static constexpr auto StyleCode = style::Code + style::NoQuote;
+static constexpr auto StyleComment = style::Comment + style::NoQuote;
+static constexpr auto StyleEnum = style::Enum + style::NoQuote;
+static constexpr auto StyleError = style::Error + style::NoQuote;
+static constexpr auto StyleFunction = style::Function + style::NoQuote;
+static constexpr auto StyleInstruction = style::Instruction + style::NoQuote;
+static constexpr auto StyleKeyword = style::Keyword + style::NoQuote;
+static constexpr auto StyleLabel = style::Label + style::NoQuote;
+static constexpr auto StyleLiteral = style::Literal + style::NoQuote;
+static constexpr auto StyleType = style::Type + style::NoQuote;
+static constexpr auto StyleVariable = style::Variable + style::NoQuote;
+
 class ScopedIndent {
   public:
     explicit ScopedIndent(uint32_t& indent) : indent_(indent) { indent_ += 2; }
@@ -87,7 +104,7 @@
 
 }  // namespace
 
-std::string Disassemble(const Module& mod) {
+StyledText Disassemble(const Module& mod) {
     return Disassembler{mod}.Disassemble();
 }
 
@@ -95,7 +112,7 @@
 
 Disassembler::~Disassembler() = default;
 
-StringStream& Disassembler::Indent() {
+StyledText& Disassembler::Indent() {
     for (uint32_t i = 0; i < indent_size_; i++) {
         out_ << " ";
     }
@@ -103,9 +120,9 @@
 }
 
 void Disassembler::EmitLine() {
-    out_ << std::endl;
+    out_ << "\n";
     current_output_line_ += 1;
-    current_output_start_pos_ = out_.tellp();
+    current_output_start_pos_ = static_cast<uint32_t>(out_.Length());
 }
 
 size_t Disassembler::IdOf(const Block* node) {
@@ -159,10 +176,17 @@
 }
 
 Source::Location Disassembler::MakeCurrentLocation() {
-    return Source::Location{current_output_line_, out_.tellp() - current_output_start_pos_ + 1};
+    return Source::Location{
+        current_output_line_,
+        static_cast<uint32_t>(out_.Length()) - current_output_start_pos_ + 1,
+    };
 }
 
-std::string Disassembler::Disassemble() {
+const StyledText& Disassembler::Disassemble() {
+    TINT_DEFER(out_ << StylePlain);
+    out_.Clear();
+    out_ << StyleCode;
+
     for (auto* ty : mod_.Types()) {
         if (auto* str = ty->As<core::type::Struct>()) {
             EmitStructDecl(str);
@@ -177,14 +201,14 @@
     for (auto& func : mod_.functions) {
         EmitFunction(func);
     }
-    return out_.str();
+    return out_;
 }
 
 void Disassembler::EmitBlock(const Block* blk, std::string_view comment /* = "" */) {
     Indent();
 
     SourceMarker sm(this);
-    out_ << "$B" << IdOf(blk);
+    out_ << StyleLabel("$B", IdOf(blk));
     if (auto* merge = blk->As<MultiInBlock>()) {
         if (!merge->Params().IsEmpty()) {
             out_ << " (";
@@ -203,7 +227,7 @@
 
     out_ << ": {";
     if (!comment.empty()) {
-        out_ << "  # " << comment;
+        out_ << "  " << StyleComment("# ", comment);
     }
 
     EmitLine();
@@ -220,17 +244,18 @@
 }
 
 void Disassembler::EmitBindingPoint(BindingPoint p) {
-    out_ << "@binding_point(" << p.group << ", " << p.binding << ")";
+    out_ << StyleAttribute("@binding_point") << "(" << StyleLiteral(p.group) << ", "
+         << StyleLiteral(p.binding) << ")";
 }
 
 void Disassembler::EmitLocation(Location loc) {
-    out_ << "@location(" << loc.value << ")";
+    out_ << StyleAttribute("@location") << "(" << loc.value << ")";
     if (loc.interpolation.has_value()) {
-        out_ << ", @interpolate(";
-        out_ << loc.interpolation->type;
+        out_ << ", " << StyleAttribute("@interpolate") << "(";
+        out_ << StyleEnum(loc.interpolation->type);
         if (loc.interpolation->sampling != core::InterpolationSampling::kUndefined) {
             out_ << ", ";
-            out_ << loc.interpolation->sampling;
+            out_ << StyleEnum(loc.interpolation->sampling);
         }
         out_ << ")";
     }
@@ -253,7 +278,7 @@
 
     if (p->Invariant()) {
         comma();
-        out_ << "@invariant";
+        out_ << StyleAttribute("@invariant");
         need_comma = true;
     }
     if (p->Location().has_value()) {
@@ -267,7 +292,7 @@
     }
     if (p->Builtin().has_value()) {
         comma();
-        out_ << "@" << p->Builtin().value();
+        out_ << StyleAttribute("@", p->Builtin().value());
         need_comma = true;
     }
     out_ << "]";
@@ -289,7 +314,7 @@
     };
     if (func->ReturnInvariant()) {
         comma();
-        out_ << "@invariant";
+        out_ << StyleAttribute("@invariant");
         need_comma = true;
     }
     if (func->ReturnLocation().has_value()) {
@@ -299,7 +324,7 @@
     }
     if (func->ReturnBuiltin().has_value()) {
         comma();
-        out_ << "@" << func->ReturnBuiltin().value();
+        out_ << StyleAttribute("@", func->ReturnBuiltin().value());
         need_comma = true;
     }
     out_ << "]";
@@ -311,32 +336,33 @@
     std::string fn_id = IdOf(func);
     {
         SourceMarker sm(this);
-        Indent() << "%" << fn_id;
+        Indent() << StyleFunction("%", fn_id);
         sm.Store(func);
     }
     out_ << " =";
 
     if (func->Stage() != Function::PipelineStage::kUndefined) {
-        out_ << " @" << func->Stage();
+        out_ << " " << StyleAttribute("@", func->Stage());
     }
     if (func->WorkgroupSize()) {
         auto arr = func->WorkgroupSize().value();
-        out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
+        out_ << " " << StyleAttribute("@workgroup_size") << "(" << StyleLiteral(arr[0]) << ", "
+             << StyleLiteral(arr[1]) << ", " << StyleLiteral(arr[2]) << ")";
     }
 
-    out_ << " func(";
+    out_ << " " << StyleKeyword("func") << "(";
 
     for (auto* p : func->Params()) {
         if (p != func->Params().Front()) {
             out_ << ", ";
         }
         SourceMarker sm(this);
-        out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+        out_ << StyleVariable("%", IdOf(p)) << ":" << StyleType(p->Type()->FriendlyName());
         sm.Store(p);
 
         EmitParamAttributes(p);
     }
-    out_ << "):" << func->ReturnType()->FriendlyName();
+    out_ << "):" << StyleType(func->ReturnType()->FriendlyName());
 
     EmitReturnAttributes(func);
 
@@ -358,7 +384,7 @@
             }
         }
         if (!names.IsEmpty()) {
-            out_ << "  # " << tint::Join(names, ", ");
+            out_ << "  " << StyleComment("# ", tint::Join(names, ", "));
         }
     }
 
@@ -389,7 +415,7 @@
     }
 
     EmitValue(val);
-    out_ << ":" << val->Type()->FriendlyName();
+    out_ << ":" << StyleType(val->Type()->FriendlyName());
 }
 
 void Disassembler::EmitValue(const Value* val) {
@@ -401,33 +427,33 @@
                     tint::Switch(
                         c,
                         [&](const core::constant::Scalar<AFloat>* scalar) {
-                            out_ << scalar->ValueAs<AFloat>().value;
+                            out_ << StyleLiteral(scalar->ValueAs<AFloat>().value);
                         },
                         [&](const core::constant::Scalar<AInt>* scalar) {
-                            out_ << scalar->ValueAs<AInt>().value;
+                            out_ << StyleLiteral(scalar->ValueAs<AInt>().value);
                         },
                         [&](const core::constant::Scalar<i32>* scalar) {
-                            out_ << scalar->ValueAs<i32>().value << "i";
+                            out_ << StyleLiteral(scalar->ValueAs<i32>().value, "i");
                         },
                         [&](const core::constant::Scalar<u32>* scalar) {
-                            out_ << scalar->ValueAs<u32>().value << "u";
+                            out_ << StyleLiteral(scalar->ValueAs<u32>().value, "u");
                         },
                         [&](const core::constant::Scalar<f32>* scalar) {
-                            out_ << scalar->ValueAs<f32>().value << "f";
+                            out_ << StyleLiteral(scalar->ValueAs<f32>().value, "f");
                         },
                         [&](const core::constant::Scalar<f16>* scalar) {
-                            out_ << scalar->ValueAs<f16>().value << "h";
+                            out_ << StyleLiteral(scalar->ValueAs<f16>().value, "h");
                         },
                         [&](const core::constant::Scalar<bool>* scalar) {
-                            out_ << (scalar->ValueAs<bool>() ? "true" : "false");
+                            out_ << StyleLiteral((scalar->ValueAs<bool>() ? "true" : "false"));
                         },
                         [&](const core::constant::Splat* splat) {
-                            out_ << splat->Type()->FriendlyName() << "(";
+                            out_ << StyleType(splat->Type()->FriendlyName()) << "(";
                             emit(splat->Index(0));
                             out_ << ")";
                         },
                         [&](const core::constant::Composite* composite) {
-                            out_ << composite->Type()->FriendlyName() << "(";
+                            out_ << StyleType(composite->Type()->FriendlyName()) << "(";
                             bool need_comma = false;
                             for (const auto* elem : composite->elements) {
                                 if (need_comma) {
@@ -441,24 +467,24 @@
                 };
             emit(constant->Value());
         },
-        [&](const ir::InstructionResult* rv) { out_ << "%" << IdOf(rv); },
+        [&](const ir::InstructionResult* rv) { out_ << StyleVariable("%", IdOf(rv)); },
         [&](const ir::BlockParam* p) {
-            out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+            out_ << StyleVariable("%", IdOf(p)) << ":" << StyleType(p->Type()->FriendlyName());
         },
-        [&](const ir::FunctionParam* p) { out_ << "%" << IdOf(p); },
-        [&](const ir::Function* f) { out_ << "%" << IdOf(f); },
+        [&](const ir::FunctionParam* p) { out_ << StyleVariable("%", IdOf(p)); },
+        [&](const ir::Function* f) { out_ << StyleVariable("%", IdOf(f)); },
         [&](Default) {
             if (val == nullptr) {
-                out_ << "undef";
+                out_ << StyleVariable("undef");
             } else {
-                out_ << "Unknown value: " << val->TypeInfo().name;
+                out_ << StyleError("unknown value: ", val->TypeInfo().name);
             }
         });
 }
 
 void Disassembler::EmitInstructionName(const Instruction* inst) {
     SourceMarker sm(this);
-    out_ << inst->FriendlyName();
+    out_ << StyleInstruction(inst->FriendlyName());
     sm.Store(inst);
 }
 
@@ -467,7 +493,7 @@
 
     if (!inst->Alive()) {
         SourceMarker sm(this);
-        out_ << "<destroyed " << inst->TypeInfo().name << " " << tint::ToString(inst) << ">";
+        out_ << StyleError("<destroyed ", inst->TypeInfo().name, " ", inst, ">");
         sm.Store(inst);
         return;
     }
@@ -515,24 +541,27 @@
                 EmitBindingPoint(v->BindingPoint().value());
             }
             if (v->Attributes().invariant) {
-                out_ << " @invariant";
+                out_ << " " << StyleAttribute("@invariant");
             }
             if (v->Attributes().location.has_value()) {
-                out_ << " @location(" << v->Attributes().location.value() << ")";
+                out_ << " " << StyleAttribute("@location") << "("
+                     << v->Attributes().location.value() << ")";
             }
             if (v->Attributes().blend_src.has_value()) {
-                out_ << " @blend_src(" << v->Attributes().blend_src.value() << ")";
+                out_ << " " << StyleAttribute("@blend_src") << "("
+                     << v->Attributes().blend_src.value() << ")";
             }
             if (v->Attributes().interpolation.has_value()) {
                 auto& interp = v->Attributes().interpolation.value();
-                out_ << " @interpolate(" << interp.type;
+                out_ << " " << StyleAttribute("@interpolate") << "(" << interp.type;
                 if (interp.sampling != core::InterpolationSampling::kUndefined) {
                     out_ << ", " << interp.sampling;
                 }
                 out_ << ")";
             }
             if (v->Attributes().builtin.has_value()) {
-                out_ << " @builtin(" << v->Attributes().builtin.value() << ")";
+                out_ << " " << StyleAttribute("@builtin") << "(" << v->Attributes().builtin.value()
+                     << ")";
             }
         },
         [&](const Swizzle* s) {
@@ -616,19 +645,19 @@
         }
         out_ << " = ";
     }
-    out_ << "if ";
+    out_ << StyleInstruction("if") << " ";
     EmitOperand(if_, If::kConditionOperandOffset);
 
     bool has_false = !if_->False()->IsEmpty();
 
-    out_ << " [t: $B" << IdOf(if_->True());
+    out_ << " [" << StyleKeyword("t") << ": " << StyleLabel("$B", IdOf(if_->True()));
     if (has_false) {
-        out_ << ", f: $B" << IdOf(if_->False());
+        out_ << ", " << StyleKeyword("f") << ": " << StyleLabel("$B", IdOf(if_->False()));
     }
     out_ << "]";
     sm.Store(if_);
 
-    out_ << " {  # " << NameOf(if_);
+    out_ << " {  " << StyleComment("# ", NameOf(if_));
     EmitLine();
 
     // True block is assumed to have instructions
@@ -643,9 +672,9 @@
     } else if (auto results = if_->Results(); !results.IsEmpty()) {
         ScopedIndent si(indent_size_);
         Indent();
-        out_ << "# implicit false block: exit_if undef";
+        out_ << StyleComment("# implicit false block: exit_if undef");
         for (size_t v = 1; v < if_->Results().Length(); v++) {
-            out_ << ", undef";
+            out_ << StyleComment(", undef");
         }
         EmitLine();
     }
@@ -655,15 +684,6 @@
 }
 
 void Disassembler::EmitLoop(const Loop* l) {
-    Vector<std::string, 3> parts;
-    if (!l->Initializer()->IsEmpty()) {
-        parts.Push("i: $B" + std::to_string(IdOf(l->Initializer())));
-    }
-    parts.Push("b: $B" + std::to_string(IdOf(l->Body())));
-
-    if (!l->Continuing()->IsEmpty()) {
-        parts.Push("c: $B" + std::to_string(IdOf(l->Continuing())));
-    }
     SourceMarker sm(this);
     if (auto results = l->Results(); !results.IsEmpty()) {
         for (size_t i = 0; i < results.Length(); ++i) {
@@ -676,10 +696,26 @@
         }
         out_ << " = ";
     }
-    out_ << "loop [" << tint::Join(parts, ", ") << "]";
+    out_ << StyleInstruction("loop") << " [";
+
+    if (!l->Initializer()->IsEmpty()) {
+        out_ << StyleKeyword("i") << ": "
+             << StyleLabel("$B", std::to_string(IdOf(l->Initializer())));
+        out_ << ", ";
+    }
+
+    out_ << StyleKeyword("b") << ": " << StyleLabel("$B", std::to_string(IdOf(l->Body())));
+
+    if (!l->Continuing()->IsEmpty()) {
+        out_ << ", ";
+        out_ << StyleKeyword("c") << ": "
+             << StyleLabel("$B", std::to_string(IdOf(l->Continuing())));
+    }
+
+    out_ << "]";
     sm.Store(l);
 
-    out_ << " {  # " << NameOf(l);
+    out_ << " {  " << StyleComment("# ", NameOf(l));
     EmitLine();
 
     if (!l->Initializer()->IsEmpty()) {
@@ -715,7 +751,7 @@
         }
         out_ << " = ";
     }
-    out_ << "switch ";
+    out_ << StyleInstruction("switch") << " ";
     EmitValue(s->Condition());
     out_ << " [";
     for (auto& c : s->Cases()) {
@@ -729,17 +765,17 @@
             }
 
             if (selector.IsDefault()) {
-                out_ << "default";
+                out_ << StyleKeyword("default");
             } else {
                 EmitValue(selector.val);
             }
         }
-        out_ << ", $B" << IdOf(c.block) << ")";
+        out_ << ", " << StyleLabel("$B", IdOf(c.block)) << ")";
     }
     out_ << "]";
     sm.Store(s);
 
-    out_ << " {  # " << NameOf(s);
+    out_ << " {  " << StyleComment("# ", NameOf(s));
     EmitLine();
 
     for (auto& c : s->Cases()) {
@@ -757,37 +793,37 @@
     tint::Switch(
         b,
         [&](const ir::Return*) {
-            out_ << "ret";
+            out_ << StyleInstruction("ret");
             args_offset = ir::Return::kArgsOperandOffset;
         },
         [&](const ir::Continue*) {
-            out_ << "continue";
+            out_ << StyleInstruction("continue");
             args_offset = ir::Continue::kArgsOperandOffset;
         },
         [&](const ir::ExitIf*) {
-            out_ << "exit_if";
+            out_ << StyleInstruction("exit_if");
             args_offset = ir::ExitIf::kArgsOperandOffset;
         },
         [&](const ir::ExitSwitch*) {
-            out_ << "exit_switch";
+            out_ << StyleInstruction("exit_switch");
             args_offset = ir::ExitSwitch::kArgsOperandOffset;
         },
         [&](const ir::ExitLoop*) {
-            out_ << "exit_loop";
+            out_ << StyleInstruction("exit_loop");
             args_offset = ir::ExitLoop::kArgsOperandOffset;
         },
         [&](const ir::NextIteration*) {
-            out_ << "next_iteration";
+            out_ << StyleInstruction("next_iteration");
             args_offset = ir::NextIteration::kArgsOperandOffset;
         },
-        [&](const ir::Unreachable*) { out_ << "unreachable"; },
+        [&](const ir::Unreachable*) { out_ << StyleInstruction("unreachable"); },
         [&](const ir::BreakIf* bi) {
-            out_ << "break_if ";
+            out_ << StyleInstruction("break_if") << " ";
             EmitValue(bi->Condition());
             args_offset = ir::BreakIf::kArgsOperandOffset;
         },
-        [&](const ir::TerminateInvocation*) { out_ << "terminate_invocation"; },
-        [&](Default) { out_ << "unknown terminator " << b->TypeInfo().name; });
+        [&](const ir::TerminateInvocation*) { out_ << StyleInstruction("terminate_invocation"); },
+        [&](Default) { out_ << StyleError("unknown terminator ", b->TypeInfo().name); });
 
     if (!b->Args().IsEmpty()) {
         out_ << " ";
@@ -798,15 +834,21 @@
     tint::Switch(
         b,  //
         [&](const ir::BreakIf* bi) {
-            out_ << "  # -> [t: exit_loop " << NameOf(bi->Loop()) << ", f: $B"
-                 << IdOf(bi->Loop()->Body()) << "]";
-        },                                                                                     //
-        [&](const ir::Continue* c) { out_ << "  # -> $B" << IdOf(c->Loop()->Continuing()); },  //
-        [&](const ir::ExitIf* e) { out_ << "  # " << NameOf(e->If()); },                       //
-        [&](const ir::ExitSwitch* e) { out_ << "  # " << NameOf(e->Switch()); },               //
-        [&](const ir::ExitLoop* e) { out_ << "  # " << NameOf(e->Loop()); },                   //
-        [&](const ir::NextIteration* ni) { out_ << "  # -> $B" << IdOf(ni->Loop()->Body()); }  //
-    );
+            out_ << "  "
+                 << StyleComment("# -> [t: exit_loop ", NameOf(bi->Loop()), ", f: $B",
+                                 IdOf(bi->Loop()->Body()), "]");
+        },
+        [&](const ir::Continue* c) {
+            out_ << "  " << StyleComment("# -> $B", IdOf(c->Loop()->Continuing()));
+        },                                                                                  //
+        [&](const ir::ExitIf* e) { out_ << "  " << StyleComment("# ", NameOf(e->If())); },  //
+        [&](const ir::ExitSwitch* e) {
+            out_ << "  " << StyleComment("# ", NameOf(e->Switch()));
+        },                                                                                      //
+        [&](const ir::ExitLoop* e) { out_ << "  " << StyleComment("# ", NameOf(e->Loop())); },  //
+        [&](const ir::NextIteration* ni) {
+            out_ << "  " << StyleComment("# -> $B", IdOf(ni->Loop()->Body()));
+        });
 }
 
 void Disassembler::EmitBinary(const Binary* b) {
@@ -815,58 +857,58 @@
     out_ << " = ";
     switch (b->Op()) {
         case BinaryOp::kAdd:
-            out_ << "add";
+            out_ << StyleInstruction("add");
             break;
         case BinaryOp::kSubtract:
-            out_ << "sub";
+            out_ << StyleInstruction("sub");
             break;
         case BinaryOp::kMultiply:
-            out_ << "mul";
+            out_ << StyleInstruction("mul");
             break;
         case BinaryOp::kDivide:
-            out_ << "div";
+            out_ << StyleInstruction("div");
             break;
         case BinaryOp::kModulo:
-            out_ << "mod";
+            out_ << StyleInstruction("mod");
             break;
         case BinaryOp::kAnd:
-            out_ << "and";
+            out_ << StyleInstruction("and");
             break;
         case BinaryOp::kOr:
-            out_ << "or";
+            out_ << StyleInstruction("or");
             break;
         case BinaryOp::kXor:
-            out_ << "xor";
+            out_ << StyleInstruction("xor");
             break;
         case BinaryOp::kEqual:
-            out_ << "eq";
+            out_ << StyleInstruction("eq");
             break;
         case BinaryOp::kNotEqual:
-            out_ << "neq";
+            out_ << StyleInstruction("neq");
             break;
         case BinaryOp::kLessThan:
-            out_ << "lt";
+            out_ << StyleInstruction("lt");
             break;
         case BinaryOp::kGreaterThan:
-            out_ << "gt";
+            out_ << StyleInstruction("gt");
             break;
         case BinaryOp::kLessThanEqual:
-            out_ << "lte";
+            out_ << StyleInstruction("lte");
             break;
         case BinaryOp::kGreaterThanEqual:
-            out_ << "gte";
+            out_ << StyleInstruction("gte");
             break;
         case BinaryOp::kShiftLeft:
-            out_ << "shl";
+            out_ << StyleInstruction("shl");
             break;
         case BinaryOp::kShiftRight:
-            out_ << "shr";
+            out_ << StyleInstruction("shr");
             break;
         case BinaryOp::kLogicalAnd:
-            out_ << "logical-and";
+            out_ << StyleInstruction("logical-and");
             break;
         case BinaryOp::kLogicalOr:
-            out_ << "logical-or";
+            out_ << StyleInstruction("logical-or");
             break;
     }
     out_ << " ";
@@ -881,19 +923,19 @@
     out_ << " = ";
     switch (u->Op()) {
         case UnaryOp::kComplement:
-            out_ << "complement";
+            out_ << StyleInstruction("complement");
             break;
         case UnaryOp::kNegation:
-            out_ << "negation";
+            out_ << StyleInstruction("negation");
             break;
         case UnaryOp::kAddressOf:
-            out_ << "ref-to-ptr";
+            out_ << StyleInstruction("ref-to-ptr");
             break;
         case UnaryOp::kIndirection:
-            out_ << "ptr-to-ref";
+            out_ << StyleInstruction("ptr-to-ref");
             break;
         case UnaryOp::kNot:
-            out_ << "not";
+            out_ << StyleInstruction("not");
             break;
     }
     out_ << " ";
@@ -903,31 +945,35 @@
 }
 
 void Disassembler::EmitStructDecl(const core::type::Struct* str) {
-    out_ << str->Name().Name() << " = struct @align(" << str->Align() << ")";
+    out_ << StyleType(str->Name().Name()) << " = " << StyleKeyword("struct") << " "
+         << StyleAttribute("@align") << "(" << StyleLiteral(str->Align()) << ")";
     if (str->StructFlags().Contains(core::type::StructFlag::kBlock)) {
-        out_ << ", @block";
+        out_ << ", " << StyleAttribute("@block");
     }
     out_ << " {";
     EmitLine();
     for (auto* member : str->Members()) {
-        out_ << "  " << member->Name().Name() << ":" << member->Type()->FriendlyName();
-        out_ << " @offset(" << member->Offset() << ")";
+        out_ << "  " << StyleVariable(member->Name().Name()) << ":"
+             << StyleType(member->Type()->FriendlyName());
+        out_ << " " << StyleAttribute("@offset") << "(" << StyleLiteral(member->Offset()) << ")";
         if (member->Attributes().invariant) {
-            out_ << ", @invariant";
+            out_ << ", " << StyleAttribute("@invariant");
         }
         if (member->Attributes().location.has_value()) {
-            out_ << ", @location(" << member->Attributes().location.value() << ")";
+            out_ << ", " << StyleAttribute("@location") << "("
+                 << StyleLiteral(member->Attributes().location.value()) << ")";
         }
         if (member->Attributes().interpolation.has_value()) {
             auto& interp = member->Attributes().interpolation.value();
-            out_ << ", @interpolate(" << interp.type;
+            out_ << ", " << StyleAttribute("@interpolate") << "(" << StyleEnum(interp.type);
             if (interp.sampling != core::InterpolationSampling::kUndefined) {
-                out_ << ", " << interp.sampling;
+                out_ << ", " << StyleEnum(interp.sampling);
             }
             out_ << ")";
         }
         if (member->Attributes().builtin.has_value()) {
-            out_ << ", @builtin(" << member->Attributes().builtin.value() << ")";
+            out_ << ", " << StyleAttribute("@builtin") << "("
+                 << StyleLiteral(member->Attributes().builtin.value()) << ")";
         }
         EmitLine();
     }
diff --git a/src/tint/lang/core/ir/disassembler.h b/src/tint/lang/core/ir/disassembler.h
index 6b6f7c6..df20534 100644
--- a/src/tint/lang/core/ir/disassembler.h
+++ b/src/tint/lang/core/ir/disassembler.h
@@ -33,7 +33,6 @@
 #include "src/tint/lang/core/ir/binary.h"
 #include "src/tint/lang/core/ir/block.h"
 #include "src/tint/lang/core/ir/block_param.h"
-#include "src/tint/lang/core/ir/call.h"
 #include "src/tint/lang/core/ir/if.h"
 #include "src/tint/lang/core/ir/loop.h"
 #include "src/tint/lang/core/ir/module.h"
@@ -42,6 +41,7 @@
 #include "src/tint/utils/containers/hashmap.h"
 #include "src/tint/utils/containers/hashset.h"
 #include "src/tint/utils/text/string_stream.h"
+#include "src/tint/utils/text/styled_text.h"
 
 // Forward declarations.
 namespace tint::core::type {
@@ -52,7 +52,7 @@
 
 /// @returns the disassembly for the module @p mod
 /// @param mod the module to disassemble
-std::string Disassemble(const Module& mod);
+StyledText Disassemble(const Module& mod);
 
 /// Helper class to disassemble the IR
 class Disassembler {
@@ -80,12 +80,9 @@
     explicit Disassembler(const Module& mod);
     ~Disassembler();
 
-    /// Returns the module as a string
+    /// Returns the module as a styled text string
     /// @returns the string representation of the module
-    std::string Disassemble();
-
-    /// @returns the string representation
-    std::string AsString() const { return out_.str(); }
+    const StyledText& Disassemble();
 
     /// @param inst the instruction to retrieve
     /// @returns the source for the instruction
@@ -190,7 +187,7 @@
         Source::Location begin_;
     };
 
-    StringStream& Indent();
+    StyledText& Indent();
 
     size_t IdOf(const Block* blk);
     std::string IdOf(const Value* node);
@@ -221,7 +218,7 @@
     void EmitInstructionName(const Instruction* inst);
 
     const Module& mod_;
-    StringStream out_;
+    StyledText out_;
     Hashmap<const Block*, size_t, 32> block_ids_;
     Hashmap<const Value*, std::string, 32> value_ids_;
     Hashset<std::string, 32> ids_;
diff --git a/src/tint/lang/core/ir/ice.h b/src/tint/lang/core/ir/ice.h
index 40ef1dd..6487cc3 100644
--- a/src/tint/lang/core/ir/ice.h
+++ b/src/tint/lang/core/ir/ice.h
@@ -31,6 +31,7 @@
 #include "src/tint/lang/core/ir/disassembler.h"
 
 /// Emit an ICE message with the disassembly of `mod` attached.
-#define TINT_IR_ICE(mod) TINT_ICE() << tint::core::ir::Disassembler{mod}.Disassemble() << "\n"
+#define TINT_IR_ICE(mod) \
+    TINT_ICE() << tint::core::ir::Disassembler{mod}.Disassemble().Plain() << "\n"
 
 #endif  // SRC_TINT_LANG_CORE_IR_ICE_H_
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
index 1270cf7..374d980 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
@@ -38,6 +38,7 @@
 #include "src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.h"
 #include "src/tint/lang/wgsl/writer/raise/raise.h"
 #include "src/tint/lang/wgsl/writer/writer.h"
+#include "src/tint/utils/text/styled_text.h"
 
 namespace tint::core::ir::transform {
 namespace {
@@ -80,7 +81,7 @@
             return "DirectVariableAccess failed:\n" + res.Failure().reason.Str();
         }
 
-        auto pre_raise = ir::Disassemble(module.Get());
+        auto pre_raise = ir::Disassemble(module.Get()).Plain();
 
         if (auto raise = wgsl::writer::Raise(module.Get()); raise != Success) {
             return "wgsl::writer::Raise failed:\n" + res.Failure().reason.Str();
@@ -89,15 +90,15 @@
         auto program_out = wgsl::writer::IRToProgram(module.Get(), program_options);
         if (!program_out.IsValid()) {
             return "wgsl::writer::IRToProgram() failed: \n" + program_out.Diagnostics().Str() +
-                   "\n\nIR (pre):\n" + pre_raise +                       //
-                   "\n\nIR (post):\n" + ir::Disassemble(module.Get()) +  //
+                   "\n\nIR (pre):\n" + pre_raise +                               //
+                   "\n\nIR (post):\n" + ir::Disassemble(module.Get()).Plain() +  //
                    "\n\nAST:\n" + Program::printer(program_out);
         }
 
         auto output = wgsl::writer::Generate(program_out, wgsl::writer::Options{});
         if (output != Success) {
             return "wgsl::writer::IRToProgram() failed: \n" + output.Failure().reason.Str() +
-                   "\n\nIR:\n" + ir::Disassemble(module.Get());
+                   "\n\nIR:\n" + ir::Disassemble(module.Get()).Plain();
         }
 
         return "\n" + output->wgsl;
diff --git a/src/tint/lang/core/ir/transform/helper_test.h b/src/tint/lang/core/ir/transform/helper_test.h
index 84e0f7e..ef8e375 100644
--- a/src/tint/lang/core/ir/transform/helper_test.h
+++ b/src/tint/lang/core/ir/transform/helper_test.h
@@ -62,7 +62,7 @@
     }
 
     /// @returns the transformed module as a disassembled string
-    std::string str() { return "\n" + ir::Disassemble(mod); }
+    std::string str() { return "\n" + ir::Disassemble(mod).Plain(); }
 
   protected:
     /// The test IR module.
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 24d594d..c9db3d3 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -357,7 +357,7 @@
     if (disassembly_file) {
         return;
     }
-    disassembly_file = std::make_unique<Source::File>("", dis_.Disassemble());
+    disassembly_file = std::make_unique<Source::File>("", dis_.Disassemble().Plain());
 }
 
 Result<SuccessType> Validator::Run() {
diff --git a/src/tint/lang/core/type/struct.cc b/src/tint/lang/core/type/struct.cc
index f403c1e..e3038b4 100644
--- a/src/tint/lang/core/type/struct.cc
+++ b/src/tint/lang/core/type/struct.cc
@@ -119,11 +119,11 @@
 }
 
 StyledText Struct::Layout() const {
-    static constexpr auto Code = style::CodeNoQuote;
-    static constexpr auto Comment = style::Comment + style::CodeNoQuote;
-    static constexpr auto Keyword = style::Keyword + style::CodeNoQuote;
-    static constexpr auto Type = style::Type + style::CodeNoQuote;
-    static constexpr auto Variable = style::Variable + style::CodeNoQuote;
+    static constexpr auto Code = style::Code + style::NoQuote;
+    static constexpr auto Comment = style::Comment + style::NoQuote;
+    static constexpr auto Keyword = style::Keyword + style::NoQuote;
+    static constexpr auto Type = style::Type + style::NoQuote;
+    static constexpr auto Variable = style::Variable + style::NoQuote;
     static constexpr auto Plain = style::Plain;
 
     StyledText out;
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index 1aece7f..cf510d3 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -82,7 +82,7 @@
         }
 
         // Return the disassembled IR module.
-        return core::ir::Disassemble(parsed.Get());
+        return core::ir::Disassemble(parsed.Get()).Plain();
     }
 };
 
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index 20a0ebb..cd0626b 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -64,7 +64,7 @@
         }
 
         // Return the disassembled IR module.
-        return "\n" + core::ir::Disassemble(ir.Get());
+        return "\n" + core::ir::Disassemble(ir.Get()).Plain();
     }
 };
 
diff --git a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
index 3b41233..0265906 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
@@ -50,7 +50,7 @@
     program_options.allowed_features = AllowedFeatures::Everything();
     auto dst = tint::wgsl::writer::IRToProgram(ir, program_options);
     if (!dst.IsValid()) {
-        std::cerr << "IR:\n" << core::ir::Disassemble(ir) << std::endl;
+        std::cerr << "IR:\n" << core::ir::Disassemble(ir).Plain() << std::endl;
         if (auto result = tint::wgsl::writer::Generate(dst, {}); result == Success) {
             std::cerr << "WGSL:\n" << result->wgsl << std::endl << std::endl;
         }
diff --git a/src/tint/lang/wgsl/ir_roundtrip_test.cc b/src/tint/lang/wgsl/ir_roundtrip_test.cc
index 6573c32..ac07621 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_test.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -73,14 +73,14 @@
             return result;
         }
 
-        result.ir_pre_raise = core::ir::Disassemble(ir_module.Get());
+        result.ir_pre_raise = core::ir::Disassemble(ir_module.Get()).Plain();
 
         if (auto res = tint::wgsl::writer::Raise(ir_module.Get()); res != Success) {
             result.err = res.Failure().reason.Str();
             return result;
         }
 
-        result.ir_post_raise = core::ir::Disassemble(ir_module.Get());
+        result.ir_post_raise = core::ir::Disassemble(ir_module.Get()).Plain();
 
         writer::ProgramOptions program_options;
         program_options.allowed_features = AllowedFeatures::Everything();
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
index cb5294e..0530639 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
@@ -53,7 +53,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<u32, 3>, read_write> = var
@@ -79,7 +79,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<u32, 3>, read_write> = var
@@ -106,7 +106,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<u32, 3>, read_write> = var
@@ -134,7 +134,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:vec3<u32> = let vec3<u32>(0u)
@@ -159,7 +159,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec3<u32>, read_write> = var
@@ -184,7 +184,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec3<u32>, read_write> = var
@@ -210,7 +210,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec3<u32>, read_write> = var
@@ -234,7 +234,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
@@ -260,7 +260,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
@@ -287,7 +287,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
@@ -312,7 +312,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, mat3x4<f32>, read_write> = var
@@ -340,7 +340,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(4) {
   foo:i32 @offset(0)
 }
@@ -374,7 +374,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(4) {
   foo:i32 @offset(0)
 }
@@ -409,7 +409,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(4) {
   foo:i32 @offset(0)
 }
@@ -447,7 +447,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(Inner = struct @align(4) {
   bar:f32 @offset(0)
 }
@@ -493,7 +493,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(Inner = struct @align(16) {
   b:i32 @offset(0)
   c:f32 @offset(4)
@@ -528,7 +528,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, array<u32, 4>, read_write> = var
@@ -551,7 +551,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec2<f32>, read_write> = var
@@ -576,7 +576,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec2<f32>, read_write> = var
@@ -602,7 +602,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec2<f32>, read_write> = var
@@ -626,7 +626,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec3<f32>, read_write> = var
@@ -650,7 +650,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, vec3<f32>, read_write> = var
@@ -682,7 +682,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(16) {
   a:i32 @offset(0)
   foo:vec4<f32> @offset(16)
@@ -713,7 +713,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:vec3<u32> = let vec3<u32>(0u)
@@ -736,7 +736,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:mat3x4<f32> = let mat3x4<f32>(vec4<f32>(0.0f))
@@ -763,7 +763,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(4) {
   foo:i32 @offset(0)
 }
@@ -799,7 +799,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(Inner = struct @align(4) {
   bar:f32 @offset(0)
 }
@@ -844,7 +844,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(Inner = struct @align(16) {
   b:i32 @offset(0)
   c:f32 @offset(4)
@@ -878,7 +878,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:vec2<f32> = let vec2<f32>(0.0f)
@@ -901,7 +901,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:vec3<f32> = let vec3<f32>(0.0f)
@@ -924,7 +924,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:vec3<f32> = let vec3<f32>(0.0f)
@@ -955,7 +955,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(MyStruct = struct @align(16) {
   a:i32 @offset(0)
   foo:vec4<f32> @offset(16)
@@ -988,7 +988,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %i:i32 = let 1i
@@ -1013,7 +1013,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %i:i32 = let 1i
@@ -1038,7 +1038,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %i:i32 = let 1i
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
index 79b7366..662de7b 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/binary_test.cc
@@ -46,7 +46,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -70,7 +70,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -93,7 +93,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -116,7 +116,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -140,7 +140,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, i32, read_write> = var
 }
 
@@ -163,7 +163,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -186,7 +186,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -210,7 +210,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -233,7 +233,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -257,7 +257,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -280,7 +280,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -304,7 +304,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -327,7 +327,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -351,7 +351,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, bool, read_write> = var
 }
 
@@ -374,7 +374,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -398,7 +398,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, bool, read_write> = var
 }
 
@@ -421,7 +421,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -445,7 +445,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -469,7 +469,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():bool {
   $B1: {
     ret true
   }
@@ -506,7 +506,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():bool {
   $B1: {
     ret true
   }
@@ -542,7 +542,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -566,7 +566,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -590,7 +590,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -614,7 +614,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -638,7 +638,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -662,7 +662,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -686,7 +686,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -710,7 +710,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -733,7 +733,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 0u
   }
@@ -757,7 +757,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, u32, read_write> = var
 }
 
@@ -782,7 +782,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():f32 {
   $B1: {
     ret 0.0f
   }
@@ -820,7 +820,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:bool):bool {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func(%p:bool):bool {
   $B1: {
     ret true
   }
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
index 7423f12..ab4b650 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/builtin_test.cc
@@ -46,7 +46,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, f32, read_write> = var, 1.0f
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
index eb3dac7..f8342c8 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/call_test.cc
@@ -48,7 +48,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():f32 {
   $B1: {
     ret 0.0f
   }
@@ -74,7 +74,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = @fragment func():void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test_function = @fragment func():void {
   $B1: {
     discard
     ret
@@ -91,7 +91,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:f32):void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func(%p:f32):void {
   $B1: {
     ret
   }
@@ -113,7 +113,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, i32, read_write> = var, 1i
 }
 
@@ -135,7 +135,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
 }
 
@@ -150,7 +150,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, f32, read_write> = var, 1.0f
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/function_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/function_test.cc
index 1af67b1..2ace620 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/function_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/function_test.cc
@@ -47,7 +47,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = @vertex func():vec4<f32> [@position] {
   $B1: {
     ret vec4<f32>(0.0f)
   }
@@ -62,7 +62,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = @fragment func():void {
   $B1: {
     ret
   }
@@ -77,7 +77,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test = @compute @workgroup_size(8, 4, 2) func():void {
   $B1: {
     ret
@@ -93,7 +93,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec3<f32> {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = func():vec3<f32> {
   $B1: {
     ret vec3<f32>(0.0f)
   }
@@ -108,7 +108,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():f32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = func():f32 {
   $B1: {
     if true [t: $B2, f: $B3] {  # if_1
       $B2: {  # true
@@ -132,7 +132,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = @vertex func():vec4<f32> [@position] {
   $B1: {
     ret vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f)
   }
@@ -148,7 +148,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test = @vertex func():vec4<f32> [@invariant, @position] {
   $B1: {
     ret vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f)
@@ -164,7 +164,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test = @fragment func():vec4<f32> [@location(1)] {
   $B1: {
     ret vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f)
@@ -182,7 +182,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test = @fragment func():vec4<f32> [@location(1), @interpolate(linear, centroid)] {
   $B1: {
     ret vec4<f32>(1.0f, 2.0f, 3.0f, 4.0f)
@@ -199,7 +199,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():f32 [@frag_depth] {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = @fragment func():f32 [@frag_depth] {
   $B1: {
     ret 1.0f
   }
@@ -215,7 +215,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():u32 [@sample_mask] {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test = @fragment func():u32 [@sample_mask] {
   $B1: {
     ret 1u
   }
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/let_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/let_test.cc
index 8b67280..761e584 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/let_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/let_test.cc
@@ -44,7 +44,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:i32 = let 42i
@@ -60,7 +60,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:i32 = let 3i
@@ -78,7 +78,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:i32 = let 1i
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/materialize_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/materialize_test.cc
index 7b1b1e6..f09beb8 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/materialize_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/materialize_test.cc
@@ -46,7 +46,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():f32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%test_function = func():f32 {
   $B1: {
     ret 2.0f
   }
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index 466498e..3e8e597 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -83,7 +83,7 @@
 
     EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():void {
   $B1: {
     ret
   }
@@ -104,7 +104,7 @@
 
     EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32):u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func(%a:u32):u32 {
   $B1: {
     ret %a
   }
@@ -126,7 +126,7 @@
 
     EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32, %b:i32, %c:bool):void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func(%a:u32, %b:i32, %c:bool):void {
   $B1: {
     ret
   }
@@ -154,7 +154,7 @@
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     if true [t: $B2, f: $B3] {  # if_1
@@ -182,7 +182,7 @@
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     if true [t: $B2] {  # if_1
@@ -207,7 +207,7 @@
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     if true [t: $B2, f: $B3] {  # if_1
@@ -235,7 +235,7 @@
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     if true [t: $B2, f: $B3] {  # if_1
@@ -262,7 +262,7 @@
 
     auto m = res.Move();
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     if true [t: $B2] {  # if_1
@@ -296,7 +296,7 @@
     EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2] {  # loop_1
@@ -326,7 +326,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -364,7 +364,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -391,7 +391,7 @@
     ASSERT_EQ(res, Success);
 
     auto m = res.Move();
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -425,7 +425,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -462,7 +462,7 @@
     EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2] {  # loop_1
@@ -500,7 +500,7 @@
     EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2] {  # loop_1
@@ -535,7 +535,7 @@
     EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2] {  # loop_1
@@ -574,7 +574,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -642,7 +642,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -682,7 +682,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2, c: $B3] {  # loop_1
@@ -722,7 +722,7 @@
     EXPECT_EQ(2u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [i: $B2, b: $B3, c: $B4] {  # loop_1
@@ -771,7 +771,7 @@
     EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [i: $B2, b: $B3] {  # loop_1
@@ -804,7 +804,7 @@
     EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
     EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     loop [b: $B2] {  # loop_1
@@ -849,7 +849,7 @@
     ASSERT_EQ(1u, cases[2].selectors.Length());
     EXPECT_TRUE(cases[2].selectors[0].IsDefault());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     switch 1i [c: (0i, $B2), c: (1i, $B3), c: (default, $B4)] {  # switch_1
@@ -897,7 +897,7 @@
 
     EXPECT_TRUE(cases[0].selectors[2].IsDefault());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     switch 1i [c: (0i 1i default, $B2)] {  # switch_1
@@ -928,7 +928,7 @@
     ASSERT_EQ(1u, cases[0].selectors.Length());
     EXPECT_TRUE(cases[0].selectors[0].IsDefault());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     switch 1i [c: (default, $B2)] {  # switch_1
@@ -968,7 +968,7 @@
 
     // This is 1 because the if is dead-code eliminated and the return doesn't happen.
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     switch 1i [c: (0i, $B2), c: (default, $B3)] {  # switch_1
@@ -1010,7 +1010,7 @@
     ASSERT_EQ(1u, cases[1].selectors.Length());
     EXPECT_TRUE(cases[1].selectors[0].IsDefault());
 
-    EXPECT_EQ(Disassemble(m),
+    EXPECT_EQ(Disassemble(m).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     switch 1i [c: (0i, $B2), c: (default, $B3)] {  # switch_1
@@ -1034,7 +1034,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%b = func():i32 {
   $B1: {
     ret 1i
@@ -1059,7 +1059,7 @@
     ASSERT_EQ(m, Success);
 
     EXPECT_EQ(
-        Disassemble(m.Get()),
+        Disassemble(m.Get()).Plain(),
         R"(%f = @fragment func(%a:vec4<f32> [@invariant, @position]):vec4<f32> [@location(1)] {
   $B1: {
     ret %a
@@ -1075,7 +1075,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = @fragment func(%a:f32 [@location(2)]):f32 [@location(1)] {
   $B1: {
     ret %a
@@ -1096,7 +1096,7 @@
     ASSERT_EQ(m, Success);
 
     EXPECT_EQ(
-        Disassemble(m.Get()),
+        Disassemble(m.Get()).Plain(),
         R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(linear, centroid)]):f32 [@location(1)] {
   $B1: {
     ret %a
@@ -1115,7 +1115,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(flat)]):f32 [@location(1)] {
   $B1: {
     ret %a
@@ -1138,7 +1138,7 @@
 
     EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():void {
   $B1: {
     ret
   }
@@ -1159,7 +1159,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func():void {
   $B1: {
     ret
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
index c3251f5..6803cd6 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/shadowing_test.cc
@@ -59,7 +59,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(S = struct @align(4) {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(S = struct @align(4) {
   i:i32 @offset(0)
 }
 
@@ -87,7 +87,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(S = struct @align(4) {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(S = struct @align(4) {
   i:i32 @offset(0)
 }
 
@@ -113,7 +113,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, i32, read_write> = var, 1i
 }
 
@@ -145,7 +145,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %i:ptr<private, i32, read_write> = var, 1i
 }
 
@@ -178,7 +178,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     if true [t: $B2] {  # if_1
@@ -217,7 +217,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     if true [t: $B2] {  # if_1
@@ -252,7 +252,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -298,7 +298,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -342,7 +342,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [i: $B2, b: $B3] {  # loop_1
@@ -386,7 +386,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [i: $B2, b: $B3] {  # loop_1
@@ -429,7 +429,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [i: $B2, b: $B3] {  # loop_1
@@ -476,7 +476,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [i: $B2, b: $B3] {  # loop_1
@@ -527,7 +527,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -581,7 +581,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -635,7 +635,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -685,7 +685,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     loop [b: $B2, c: $B3] {  # loop_1
@@ -735,7 +735,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     %3:i32 = load %i
@@ -783,7 +783,7 @@
 
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():i32 {
   $B1: {
     %i:ptr<function, i32, read_write> = var
     %3:i32 = load %i
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/store_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/store_test.cc
index b2ac91e..b8c04a3 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/store_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/store_test.cc
@@ -47,7 +47,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %a:ptr<private, u32, read_write> = var
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
index 8342e48..dbf7b24 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/unary_test.cc
@@ -46,7 +46,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():bool {
   $B1: {
     ret false
   }
@@ -70,7 +70,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():vec4<bool> {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():vec4<bool> {
   $B1: {
     ret vec4<bool>(false)
   }
@@ -94,7 +94,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():u32 {
   $B1: {
     ret 1u
   }
@@ -118,7 +118,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%my_func = func():i32 {
   $B1: {
     ret 1i
   }
@@ -143,7 +143,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, i32, read_write> = var
 }
 
@@ -167,7 +167,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %v1:ptr<private, i32, read_write> = var
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/var_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/var_test.cc
index c683285..16cf0c9 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/var_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/var_test.cc
@@ -45,7 +45,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %a:ptr<private, u32, read_write> = var
 }
 
@@ -59,7 +59,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %a:ptr<private, u32, read_write> = var, 2u
 }
 
@@ -72,7 +72,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()), R"($B1: {  # root
+    EXPECT_EQ(Disassemble(m.Get()).Plain(), R"($B1: {  # root
   %a:ptr<storage, u32, read> = var @binding_point(2, 3)
 }
 
@@ -86,7 +86,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, u32, read_write> = var
@@ -104,7 +104,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, u32, read_write> = var, 2u
@@ -122,7 +122,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, u32, read_write> = var
@@ -142,7 +142,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, i32, read_write> = var
@@ -179,7 +179,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func(%p:i32):i32 {
   $B1: {
     ret %p
@@ -225,7 +225,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func(%p:i32):i32 {
   $B1: {
     ret %p
@@ -273,7 +273,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func(%p:i32):i32 {
   $B1: {
     ret %p
@@ -305,7 +305,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void {
   $B1: {
     %a:ptr<function, i32, read_write> = var
@@ -344,7 +344,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func(%p:i32):i32 {
   $B1: {
     ret %p
@@ -396,7 +396,7 @@
     auto m = Build();
     ASSERT_EQ(m, Success);
 
-    EXPECT_EQ(Disassemble(m.Get()),
+    EXPECT_EQ(Disassemble(m.Get()).Plain(),
               R"(%f = func(%p:i32):i32 {
   $B1: {
     ret %p
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 30e6795..68d4db1 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -51,7 +51,7 @@
 IRToProgramTest::Result IRToProgramTest::Run() {
     Result result;
 
-    result.ir = tint::core::ir::Disassemble(mod);
+    result.ir = tint::core::ir::Disassemble(mod).Plain();
 
     ProgramOptions options;
     options.allowed_features = AllowedFeatures::Everything();
diff --git a/src/tint/lang/wgsl/writer/raise/ptr_to_ref_test.cc b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_test.cc
index b663af0..3862749 100644
--- a/src/tint/lang/wgsl/writer/raise/ptr_to_ref_test.cc
+++ b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_test.cc
@@ -67,7 +67,7 @@
     /// @returns the transformed module as a disassembled string
     std::string str() {
         core::ir::Disassembler dis(mod);
-        return "\n" + dis.Disassemble();
+        return "\n" + dis.Disassemble().Plain();
     }
 
   protected:
diff --git a/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc b/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc
index 1c805b1..7be25fc 100644
--- a/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc
+++ b/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc
@@ -67,7 +67,7 @@
     /// @returns the transformed module as a disassembled string
     std::string str() {
         core::ir::Disassembler dis(mod);
-        return "\n" + dis.Disassemble();
+        return "\n" + dis.Disassemble().Plain();
     }
 
   protected:
diff --git a/src/tint/lang/wgsl/writer/writer_test.cc b/src/tint/lang/wgsl/writer/writer_test.cc
index d01d1df..ca2f339 100644
--- a/src/tint/lang/wgsl/writer/writer_test.cc
+++ b/src/tint/lang/wgsl/writer/writer_test.cc
@@ -65,14 +65,14 @@
     Result Run(std::string_view expected_wgsl) {
         Result result;
 
-        result.ir_pre_raise = core::ir::Disassemble(mod);
+        result.ir_pre_raise = core::ir::Disassemble(mod).Plain();
 
         if (auto res = tint::wgsl::writer::Raise(mod); res != Success) {
             result.err = res.Failure().reason.Str();
             return result;
         }
 
-        result.ir_post_raise = core::ir::Disassemble(mod);
+        result.ir_post_raise = core::ir::Disassemble(mod).Plain();
 
         writer::ProgramOptions program_options;
         program_options.allowed_features = AllowedFeatures::Everything();
diff --git a/src/tint/utils/text/string_stream.h b/src/tint/utils/text/string_stream.h
index ab96301..45aaaab 100644
--- a/src/tint/utils/text/string_stream.h
+++ b/src/tint/utils/text/string_stream.h
@@ -197,8 +197,8 @@
         return *this;
     }
 
-    /// @returns the current location in the output stream
-    uint32_t tellp() { return static_cast<uint32_t>(sstream_.tellp()); }
+    /// @returns the number of UTF-8 code units (bytes) have been written to the string.
+    size_t Length() { return static_cast<size_t>(sstream_.tellp()); }
 
     /// @returns the string contents of the stream
     std::string str() const { return sstream_.str(); }
diff --git a/src/tint/utils/text/styled_text.cc b/src/tint/utils/text/styled_text.cc
index e9290ae..b6ad72e 100644
--- a/src/tint/utils/text/styled_text.cc
+++ b/src/tint/utils/text/styled_text.cc
@@ -74,10 +74,10 @@
     StringStream ss;
     bool is_code_no_quote = false;
     Walk([&](std::string_view text, TextStyle style) {
-        if (is_code_no_quote != (style.IsCode() && !style.IsCodeNoQuote())) {
+        if (is_code_no_quote != (style.IsCode() && !style.IsNoQuote())) {
             ss << "'";
         }
-        is_code_no_quote = style.IsCode() && !style.IsCodeNoQuote();
+        is_code_no_quote = style.IsCode() && !style.IsNoQuote();
 
         ss << text;
     });
diff --git a/src/tint/utils/text/styled_text.h b/src/tint/utils/text/styled_text.h
index 21a82aa..556099b 100644
--- a/src/tint/utils/text/styled_text.h
+++ b/src/tint/utils/text/styled_text.h
@@ -117,9 +117,9 @@
                        value.values);
             *this << old_style;
         } else {
-            uint32_t offset = stream_.tellp();
+            size_t offset = stream_.Length();
             stream_ << value;
-            spans_.Back().length += stream_.tellp() - offset;
+            spans_.Back().length += stream_.Length() - offset;
         }
         return *this;
     }
@@ -136,6 +136,9 @@
         }
     }
 
+    /// @returns the number of UTF-8 code units (bytes) have been written to the string.
+    size_t Length() { return stream_.Length(); }
+
   private:
     struct Span {
         TextStyle style;
diff --git a/src/tint/utils/text/styled_text_theme.cc b/src/tint/utils/text/styled_text_theme.cc
index 0785c47..0b209a1 100644
--- a/src/tint/utils/text/styled_text_theme.cc
+++ b/src/tint/utils/text/styled_text_theme.cc
@@ -135,6 +135,20 @@
         /* bold */ std::nullopt,
         /* underlined */ std::nullopt,
     },
+    /* kind_label */
+    StyledTextTheme::Attributes{
+        /* foreground */ Color{180, 140, 140},
+        /* background */ std::nullopt,
+        /* bold */ std::nullopt,
+        /* underlined */ std::nullopt,
+    },
+    /* kind_instruction */
+    StyledTextTheme::Attributes{
+        /* foreground */ Color{220, 220, 170},
+        /* background */ std::nullopt,
+        /* bold */ std::nullopt,
+        /* underlined */ std::nullopt,
+    },
     /* kind_squiggle */
     StyledTextTheme::Attributes{
         /* foreground */ Color{0, 200, 255},
@@ -249,6 +263,20 @@
         /* bold */ std::nullopt,
         /* underlined */ std::nullopt,
     },
+    /* kind_label */
+    StyledTextTheme::Attributes{
+        /* foreground */ Color{180, 140, 140},
+        /* background */ std::nullopt,
+        /* bold */ std::nullopt,
+        /* underlined */ std::nullopt,
+    },
+    /* kind_instruction */
+    StyledTextTheme::Attributes{
+        /* foreground */ Color{121, 94, 38},
+        /* background */ std::nullopt,
+        /* bold */ std::nullopt,
+        /* underlined */ std::nullopt,
+    },
     /* kind_squiggle */
     StyledTextTheme::Attributes{
         /* foreground */ Color{0, 200, 255},
@@ -310,6 +338,10 @@
                 apply(kind_attribute);
             } else if (text_style.IsComment()) {
                 apply(kind_comment);
+            } else if (text_style.IsLabel()) {
+                apply(kind_label);
+            } else if (text_style.IsInstruction()) {
+                apply(kind_instruction);
             }
         }
         if (text_style.IsSquiggle()) {
diff --git a/src/tint/utils/text/styled_text_theme.h b/src/tint/utils/text/styled_text_theme.h
index bff0de7..a184b55 100644
--- a/src/tint/utils/text/styled_text_theme.h
+++ b/src/tint/utils/text/styled_text_theme.h
@@ -105,6 +105,10 @@
     Attributes kind_attribute;
     /// The theme's attributes for a comment token. This is applied on top #kind_code.
     Attributes kind_comment;
+    /// The theme's attributes for a label token. This is applied on top #kind_code.
+    Attributes kind_label;
+    /// The theme's attributes for a instruction token. This is applied on top #kind_code.
+    Attributes kind_instruction;
 
     /// The theme's attributes for a squiggle-highlight.
     Attributes kind_squiggle;
diff --git a/src/tint/utils/text/text_style.h b/src/tint/utils/text/text_style.h
index c337a5a..2006e5d 100644
--- a/src/tint/utils/text/text_style.h
+++ b/src/tint/utils/text/text_style.h
@@ -31,12 +31,9 @@
 #define SRC_TINT_UTILS_TEXT_TEXT_STYLE_H_
 
 #include <cstdint>
-#include <string_view>
 #include <tuple>
 #include <type_traits>
 
-#include "src/tint/utils/containers/enum_set.h"
-
 // Forward declarations
 namespace tint {
 template <typename... VALUES>
@@ -53,38 +50,41 @@
 
     /// Bit patterns
 
-    static constexpr Bits kStyleMask /*          */ = 0b00'000000'0000'00'11;
-    static constexpr Bits kStyleUnderlined /*    */ = 0b00'000000'0000'00'01;
-    static constexpr Bits kStyleBold /*          */ = 0b00'000000'0000'00'10;
+    static constexpr Bits kStyleMask /*          */ = 0b00'00000'0000'00'111;
+    static constexpr Bits kStyleUnderlined /*    */ = 0b00'00000'0000'00'001;
+    static constexpr Bits kStyleBold /*          */ = 0b00'00000'0000'00'010;
+    static constexpr Bits kStyleNoQuote /*       */ = 0b00'00000'0000'00'100;
 
-    static constexpr Bits kCompareMask /*        */ = 0b00'000000'0000'11'00;
-    static constexpr Bits kCompareMatch /*       */ = 0b00'000000'0000'01'00;
-    static constexpr Bits kCompareMismatch /*    */ = 0b00'000000'0000'10'00;
+    static constexpr Bits kCompareMask /*        */ = 0b00'00000'0000'11'000;
+    static constexpr Bits kCompareMatch /*       */ = 0b00'00000'0000'01'000;
+    static constexpr Bits kCompareMismatch /*    */ = 0b00'00000'0000'10'000;
 
-    static constexpr Bits kSeverityMask /*       */ = 0b00'000000'1111'00'00;
-    static constexpr Bits kSeverityDefault /*    */ = 0b00'000000'0000'00'00;
-    static constexpr Bits kSeveritySuccess /*    */ = 0b00'000000'0001'00'00;
-    static constexpr Bits kSeverityWarning /*    */ = 0b00'000000'0010'00'00;
-    static constexpr Bits kSeverityError /*      */ = 0b00'000000'0011'00'00;
-    static constexpr Bits kSeverityFatal /*      */ = 0b00'000000'0100'00'00;
+    static constexpr Bits kSeverityMask /*       */ = 0b00'00000'1111'00'000;
+    static constexpr Bits kSeverityDefault /*    */ = 0b00'00000'0000'00'000;
+    static constexpr Bits kSeveritySuccess /*    */ = 0b00'00000'0001'00'000;
+    static constexpr Bits kSeverityWarning /*    */ = 0b00'00000'0010'00'000;
+    static constexpr Bits kSeverityError /*      */ = 0b00'00000'0011'00'000;
+    static constexpr Bits kSeverityFatal /*      */ = 0b00'00000'0100'00'000;
 
-    static constexpr Bits kKindMask /*           */ = 0b00'111111'0000'00'00;
-    static constexpr Bits kKindCode /*           */ = 0b00'000001'0000'00'00;
-    static constexpr Bits kKindCodeNoQuote /*    */ = 0b00'000011'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindKeyword /*        */ = 0b00'000101'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindVariable /*       */ = 0b00'001001'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindType /*           */ = 0b00'001101'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindFunction /*       */ = 0b00'010001'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindEnum /*           */ = 0b00'010101'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindLiteral /*        */ = 0b00'011001'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindAttribute /*      */ = 0b00'011101'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindComment /*        */ = 0b00'100001'0000'00'00;  // includes kKindCode
-    static constexpr Bits kKindSquiggle /*       */ = 0b00'000100'0000'00'00;
+    static constexpr Bits kKindMask /*           */ = 0b00'11111'0000'00'000;
+    static constexpr Bits kKindCode /*           */ = 0b00'00001'0000'00'000;
+    static constexpr Bits kKindKeyword /*        */ = 0b00'00011'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindVariable /*       */ = 0b00'00101'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindType /*           */ = 0b00'00111'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindFunction /*       */ = 0b00'01001'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindEnum /*           */ = 0b00'01011'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindLiteral /*        */ = 0b00'01101'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindAttribute /*      */ = 0b00'01111'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindComment /*        */ = 0b00'10001'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindLabel /*          */ = 0b00'10011'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindInstruction /*    */ = 0b00'10101'0000'00'000;  // includes kKindCode
+    static constexpr Bits kKindSquiggle /*       */ = 0b00'00010'0000'00'000;
 
     /// Getters
 
     constexpr bool IsBold() const { return (bits & kStyleBold) != 0; }
     constexpr bool IsUnderlined() const { return (bits & kStyleUnderlined) == kStyleUnderlined; }
+    constexpr bool IsNoQuote() const { return (bits & kStyleNoQuote) == kStyleNoQuote; }
 
     constexpr bool HasCompare() const { return (bits & kCompareMask) != 0; }
     constexpr bool IsMatch() const { return (bits & kCompareMask) == kCompareMatch; }
@@ -98,7 +98,6 @@
 
     constexpr bool HasKind() const { return (bits & kKindMask) != 0; }
     constexpr bool IsCode() const { return (bits & kKindCode) == kKindCode; }
-    constexpr bool IsCodeNoQuote() const { return (bits & kKindCodeNoQuote) == kKindCodeNoQuote; }
     constexpr bool IsKeyword() const { return (bits & kKindMask) == kKindKeyword; }
     constexpr bool IsVariable() const { return (bits & kKindMask) == kKindVariable; }
     constexpr bool IsType() const { return (bits & kKindMask) == kKindType; }
@@ -107,6 +106,8 @@
     constexpr bool IsLiteral() const { return (bits & kKindMask) == kKindLiteral; }
     constexpr bool IsAttribute() const { return (bits & kKindMask) == kKindAttribute; }
     constexpr bool IsComment() const { return (bits & kKindMask) == kKindComment; }
+    constexpr bool IsLabel() const { return (bits & kKindMask) == kKindLabel; }
+    constexpr bool IsInstruction() const { return (bits & kKindMask) == kKindInstruction; }
     constexpr bool IsSquiggle() const { return (bits & kKindMask) == kKindSquiggle; }
 
     /// Equality operator
@@ -174,9 +175,11 @@
 static constexpr TextStyle Bold = TextStyle{TextStyle::kStyleBold};
 /// Underlined renders text with an underline
 static constexpr TextStyle Underlined = TextStyle{TextStyle::kStyleUnderlined};
-/// Underlined renders text with the compare-match style
+/// Code renders without a single quote when printed as plain text.
+static constexpr TextStyle NoQuote = TextStyle{TextStyle::kStyleNoQuote};
+/// Match renders text with the compare-match style
 static constexpr TextStyle Match = TextStyle{TextStyle::kCompareMatch};
-/// Underlined renders text with the compare-mismatch style
+/// Mismatch renders text with the compare-mismatch style
 static constexpr TextStyle Mismatch = TextStyle{TextStyle::kCompareMismatch};
 /// Success renders text with the styling for a successful status
 static constexpr TextStyle Success = TextStyle{TextStyle::kSeveritySuccess};
@@ -188,9 +191,6 @@
 static constexpr TextStyle Fatal = TextStyle{TextStyle::kSeverityFatal};
 /// Code renders text with a 'code' style
 static constexpr TextStyle Code = TextStyle{TextStyle::kKindCode};
-/// Code renders text with a 'code' style, but does not use a single quote when printed as plain
-/// text.
-static constexpr TextStyle CodeNoQuote = TextStyle{TextStyle::kKindCodeNoQuote};
 /// Keyword renders text with a 'code' style that represents a 'keyword' token
 static constexpr TextStyle Keyword = TextStyle{TextStyle::kKindKeyword};
 /// Variable renders text with a 'code' style that represents a 'variable' token
@@ -207,6 +207,10 @@
 static constexpr TextStyle Attribute = TextStyle{TextStyle::kKindAttribute};
 /// Comment renders text with a 'code' style that represents an 'comment' token
 static constexpr TextStyle Comment = TextStyle{TextStyle::kKindComment};
+/// Label renders text with a 'code' style that represents an 'label' token
+static constexpr TextStyle Label = TextStyle{TextStyle::kKindLabel};
+/// Instruction renders text with a 'code' style that represents an 'instruction' token
+static constexpr TextStyle Instruction = TextStyle{TextStyle::kKindInstruction};
 /// Squiggle renders text with a squiggle-highlight style (`^^^^^`)
 static constexpr TextStyle Squiggle = TextStyle{TextStyle::kKindSquiggle};