diff --git a/src/tint/cmd/common/helper.cc b/src/tint/cmd/common/helper.cc
index c043588..2e9ce37 100644
--- a/src/tint/cmd/common/helper.cc
+++ b/src/tint/cmd/common/helper.cc
@@ -453,6 +453,8 @@
             return "Rgba32Sint";
         case tint::inspector::ResourceBinding::TexelFormat::kRgba32Float:
             return "Rgba32Float";
+        case tint::inspector::ResourceBinding::TexelFormat::kR8Unorm:
+            return "R8Unorm";
         case tint::inspector::ResourceBinding::TexelFormat::kNone:
             return "None";
     }
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index f8deb52..7352a85 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -302,7 +302,8 @@
 
 // https://gpuweb.github.io/gpuweb/wgsl/#texel-formats
 match f32_texel_format
-  : texel_format.bgra8unorm
+  : texel_format.r8unorm
+  | texel_format.bgra8unorm
   | texel_format.rgba8unorm
   | texel_format.rgba8snorm
   | texel_format.rgba16float
diff --git a/src/tint/lang/core/intrinsic/data.cc b/src/tint/lang/core/intrinsic/data.cc
index 46637b1..db57c64 100644
--- a/src/tint/lang/core/intrinsic/data.cc
+++ b/src/tint/lang/core/intrinsic/data.cc
@@ -81,7 +81,7 @@
     return BuildBool(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "bool";
+    out << style::Type("bool");
   }
 };
 
@@ -95,7 +95,7 @@
     return BuildI32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "i32";
+    out << style::Type("i32");
   }
 };
 
@@ -109,7 +109,7 @@
     return BuildU32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "u32";
+    out << style::Type("u32");
   }
 };
 
@@ -123,7 +123,7 @@
     return BuildF32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f32";
+    out << style::Type("f32");
   }
 };
 
@@ -137,7 +137,7 @@
     return BuildF16(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f16";
+    out << style::Type("f16");
   }
 };
 
@@ -157,7 +157,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec2", "<", T, ">");
   }
 };
 
@@ -177,7 +177,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec3", "<", T, ">");
   }
 };
 
@@ -197,7 +197,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec4", "<", T, ">");
   }
 };
 
@@ -217,7 +217,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x2", "<", T, ">");
   }
 };
 
@@ -237,7 +237,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x3", "<", T, ">");
   }
 };
 
@@ -257,7 +257,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x4", "<", T, ">");
   }
 };
 
@@ -277,7 +277,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x2", "<", T, ">");
   }
 };
 
@@ -297,7 +297,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x3", "<", T, ">");
   }
 };
 
@@ -317,7 +317,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x4", "<", T, ">");
   }
 };
 
@@ -337,7 +337,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x2", "<", T, ">");
   }
 };
 
@@ -357,7 +357,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x3", "<", T, ">");
   }
 };
 
@@ -377,7 +377,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x4", "<", T, ">");
   }
 };
 
@@ -403,7 +403,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "vec" << style::Type << N << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("vec", N, "<", T, ">");
   }
 };
 
@@ -435,7 +435,7 @@
   state->PrintNum(N);StyledText M;
   state->PrintNum(M);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "mat" << style::Type << N << style::Type << "x" << style::Type << M << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("mat", N, "x", M, "<", T, ">");
   }
 };
 
@@ -467,7 +467,7 @@
   state->PrintNum(S);StyledText T;
   state->PrintType(T);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "ptr" << style::Code << "<" << style::Type << S << style::Code << ", " << style::Type << T << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("ptr", "<", S, ", ", T, ", ", A, ">");
   }
 };
 
@@ -487,7 +487,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "atomic" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("atomic", "<", T, ">");
   }
 };
 
@@ -507,7 +507,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("array", "<", T, ">");
   }
 };
 
@@ -521,7 +521,7 @@
     return BuildSampler(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler";
+    out << style::Type("sampler");
   }
 };
 
@@ -535,7 +535,7 @@
     return BuildSamplerComparison(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler_comparison";
+    out << style::Type("sampler_comparison");
   }
 };
 
@@ -555,7 +555,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_1d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_1d", "<", T, ">");
   }
 };
 
@@ -575,7 +575,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d", "<", T, ">");
   }
 };
 
@@ -595,7 +595,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d_array", "<", T, ">");
   }
 };
 
@@ -615,7 +615,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_3d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_3d", "<", T, ">");
   }
 };
 
@@ -635,7 +635,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube", "<", T, ">");
   }
 };
 
@@ -655,7 +655,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube_array", "<", T, ">");
   }
 };
 
@@ -675,7 +675,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_multisampled_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_multisampled_2d", "<", T, ">");
   }
 };
 
@@ -689,7 +689,7 @@
     return BuildTextureDepth2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d";
+    out << style::Type("texture_depth_2d");
   }
 };
 
@@ -703,7 +703,7 @@
     return BuildTextureDepth2DArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d_array";
+    out << style::Type("texture_depth_2d_array");
   }
 };
 
@@ -717,7 +717,7 @@
     return BuildTextureDepthCube(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube";
+    out << style::Type("texture_depth_cube");
   }
 };
 
@@ -731,7 +731,7 @@
     return BuildTextureDepthCubeArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube_array";
+    out << style::Type("texture_depth_cube_array");
   }
 };
 
@@ -745,7 +745,7 @@
     return BuildTextureDepthMultisampled2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_multisampled_2d";
+    out << style::Type("texture_depth_multisampled_2d");
   }
 };
 
@@ -771,7 +771,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_1d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_1d", "<", F, ", ", A, ">");
   }
 };
 
@@ -797,7 +797,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d", "<", F, ", ", A, ">");
   }
 };
 
@@ -823,7 +823,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d_array" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d_array", "<", F, ", ", A, ">");
   }
 };
 
@@ -849,7 +849,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_3d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_3d", "<", F, ", ", A, ">");
   }
 };
 
@@ -863,7 +863,7 @@
     return BuildTextureExternal(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_external";
+    out << style::Type("texture_external");
   }
 };
 
@@ -883,7 +883,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__modf_result_" << style::Type << T;
+    out << style::Type("__modf_result_", T);
   }
 };
 
@@ -909,7 +909,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__modf_result_vec" << style::Type << N << style::Type << "_" << style::Type << T;
+    out << style::Type("__modf_result_vec", N, "_", T);
   }
 };
 
@@ -929,7 +929,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__frexp_result_" << style::Type << T;
+    out << style::Type("__frexp_result_", T);
   }
 };
 
@@ -955,7 +955,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__frexp_result_vec" << style::Type << N << style::Type << "_" << style::Type << T;
+    out << style::Type("__frexp_result_vec", N, "_", T);
   }
 };
 
@@ -975,7 +975,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "__atomic_compare_exchange_result" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("__atomic_compare_exchange_result", "<", T, ">");
   }
 };
 
@@ -1003,7 +1003,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_f32'
@@ -1026,7 +1026,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kI32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_f16'
@@ -1049,7 +1049,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_i32'
@@ -1072,7 +1072,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_u32'
@@ -1095,7 +1095,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_bool'
@@ -1118,7 +1118,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fiu32_f16'
@@ -1141,7 +1141,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fiu32'
@@ -1161,7 +1161,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fi32_f16'
@@ -1181,7 +1181,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fi32'
@@ -1198,7 +1198,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kI32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kI32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match f32_f16'
@@ -1215,7 +1215,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match iu32'
@@ -1232,13 +1232,14 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// EnumMatcher for 'match f32_texel_format'
 constexpr NumberMatcher kF32TexelFormatMatcher {
 /* match */ [](MatchState&, Number number) -> Number {
     switch (static_cast<core::TexelFormat>(number.Value())) {
+      case core::TexelFormat::kR8Unorm:
       case core::TexelFormat::kBgra8Unorm:
       case core::TexelFormat::kRgba8Unorm:
       case core::TexelFormat::kRgba8Snorm:
@@ -1252,7 +1253,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "bgra8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8snorm"<< TextStyle{} << ", " << style::Enum << "rgba16float"<< TextStyle{} << ", " << style::Enum << "r32float"<< TextStyle{} << ", " << style::Enum << "rg32float"<< TextStyle{} << " or " << style::Enum << "rgba32float";
+  out<< style::Enum("r8unorm")<< style::Plain(", ") << style::Enum("bgra8unorm")<< style::Plain(", ") << style::Enum("rgba8unorm")<< style::Plain(", ") << style::Enum("rgba8snorm")<< style::Plain(", ") << style::Enum("rgba16float")<< style::Plain(", ") << style::Enum("r32float")<< style::Plain(", ") << style::Enum("rg32float")<< style::Plain(" or ") << style::Enum("rgba32float");
   }
 };
 
@@ -1271,7 +1272,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8sint"<< TextStyle{} << ", " << style::Enum << "rgba16sint"<< TextStyle{} << ", " << style::Enum << "r32sint"<< TextStyle{} << ", " << style::Enum << "rg32sint"<< TextStyle{} << " or " << style::Enum << "rgba32sint";
+  out<< style::Enum("rgba8sint")<< style::Plain(", ") << style::Enum("rgba16sint")<< style::Plain(", ") << style::Enum("r32sint")<< style::Plain(", ") << style::Enum("rg32sint")<< style::Plain(" or ") << style::Enum("rgba32sint");
   }
 };
 
@@ -1290,7 +1291,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8uint"<< TextStyle{} << ", " << style::Enum << "rgba16uint"<< TextStyle{} << ", " << style::Enum << "r32uint"<< TextStyle{} << ", " << style::Enum << "rg32uint"<< TextStyle{} << " or " << style::Enum << "rgba32uint";
+  out<< style::Enum("rgba8uint")<< style::Plain(", ") << style::Enum("rgba16uint")<< style::Plain(", ") << style::Enum("r32uint")<< style::Plain(", ") << style::Enum("rg32uint")<< style::Plain(" or ") << style::Enum("rgba32uint");
   }
 };
 
@@ -1303,7 +1304,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "write";
+  out<< style::Enum("write");
   }
 };
 
@@ -1316,7 +1317,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read_write";
+  out<< style::Enum("read_write");
   }
 };
 
@@ -1332,7 +1333,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("read")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
@@ -1348,7 +1349,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "write"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("write")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
@@ -1365,7 +1366,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "function"<< TextStyle{} << ", " << style::Enum << "private"<< TextStyle{} << " or " << style::Enum << "workgroup";
+  out<< style::Enum("function")<< style::Plain(", ") << style::Enum("private")<< style::Plain(" or ") << style::Enum("workgroup");
   }
 };
 
@@ -1381,7 +1382,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "workgroup"<< TextStyle{} << " or " << style::Enum << "storage";
+  out<< style::Enum("workgroup")<< style::Plain(" or ") << style::Enum("storage");
   }
 };
 
@@ -1394,7 +1395,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "storage";
+  out<< style::Enum("storage");
   }
 };
 
@@ -1407,7 +1408,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "workgroup";
+  out<< style::Enum("workgroup");
   }
 };
 
diff --git a/src/tint/lang/core/intrinsic/table.cc b/src/tint/lang/core/intrinsic/table.cc
index b888362..1a9dae9 100644
--- a/src/tint/lang/core/intrinsic/table.cc
+++ b/src/tint/lang/core/intrinsic/table.cc
@@ -97,17 +97,16 @@
                      [&](const Candidate& a, const Candidate& b) { return a.score < b.score; });
 }
 
-/// Prints a list of types. Ends with the style of @p ss set to style::Code.
+/// Prints a list of types.
 static void PrintTypeList(StyledText& ss, VectorRef<const core::type::Type*> types) {
     bool first = true;
     for (auto* arg : types) {
         if (!first) {
-            ss << style::Code << ", ";
+            ss << ", ";
         }
         first = false;
-        ss << style::Type << arg->FriendlyName();
+        ss << style::Type(arg->FriendlyName());
     }
-    ss << style::Code;
 }
 
 /// Attempts to find a single intrinsic overload that matches the provided argument types.
@@ -191,13 +190,12 @@
                                 VectorRef<const core::type::Type*> args,
                                 VectorRef<Candidate> candidates);
 
-/// @return a string representing a call to a builtin with the given argument
-/// types.
+/// @return a string representing a call to a builtin with the given argument types.
 StyledText CallSignature(std::string_view intrinsic_name,
                          VectorRef<const core::type::Type*> template_args,
                          VectorRef<const core::type::Type*> args) {
     StyledText out;
-    out << style::Function << intrinsic_name << style::Code;
+    out << style::Code << style::Function(intrinsic_name);
     if (!template_args.IsEmpty()) {
         out << "<";
         PrintTypeList(out, template_args);
@@ -207,7 +205,7 @@
     PrintTypeList(out, args);
     out << ")";
 
-    return out << style::Plain;
+    return out;
 }
 
 Result<Overload, StyledText> MatchIntrinsic(Context& context,
@@ -541,22 +539,24 @@
                    Context& context,
                    const OverloadInfo& overload,
                    std::string_view intrinsic_name) {
-    TINT_DEFER(ss << style::Plain);
+    // Restore old style before returning.
+    auto prev_style = ss.Style();
+    TINT_DEFER(ss << prev_style);
 
     TemplateState templates;
 
     // TODO(crbug.com/tint/1730): Use input evaluation stage to output only relevant overloads.
     auto earliest_eval_stage = EvaluationStage::kConstant;
 
-    ss << style::Function << intrinsic_name << style::Code;
+    ss << style::Code << style::Function(intrinsic_name);
 
     if (overload.num_explicit_templates > 0) {
         ss << "<";
         for (size_t i = 0; i < overload.num_explicit_templates; i++) {
             if (i > 0) {
-                ss << style::Code << ", ";
+                ss << ", ";
             }
-            ss << style::Type << context.data[overload.templates + i].name;
+            ss << style::Type(context.data[overload.templates + i].name);
         }
         ss << ">";
     }
@@ -564,15 +564,15 @@
     for (size_t p = 0; p < overload.num_parameters; p++) {
         auto& parameter = context.data[overload.parameters + p];
         if (p > 0) {
-            ss << style::Code << ", ";
+            ss << ", ";
         }
         if (parameter.usage != ParameterUsage::kNone) {
-            ss << style::Variable << ToString(parameter.usage) << style::Code << ": ";
+            ss << style::Variable(parameter.usage, ": ");
         }
         auto* matcher_indices = context.data[parameter.matcher_indices];
         Match(context, templates, overload, matcher_indices, earliest_eval_stage).PrintType(ss);
     }
-    ss << style::Code << ")";
+    ss << ")";
     if (overload.return_matcher_indices.IsValid()) {
         ss << " -> ";
         auto* matcher_indices = context.data[overload.return_matcher_indices];
@@ -581,7 +581,7 @@
 
     bool first = true;
     auto separator = [&] {
-        ss << style::Plain << (first ? "  where: " : ", ");
+        ss << style::Plain(first ? "  where: " : ", ");
         first = false;
     };
 
@@ -592,7 +592,7 @@
                 Match(context, templates, overload, matcher_indices, earliest_eval_stage);
 
             separator();
-            ss << style::Type << tmpl.name << style::Plain << " is ";
+            ss << style::Type(tmpl.name) << style::Plain(" is ");
             if (tmpl.kind == TemplateInfo::Kind::kType) {
                 matcher.PrintType(ss);
             } else {
diff --git a/src/tint/lang/core/intrinsic/table_data.h b/src/tint/lang/core/intrinsic/table_data.h
index 1ed619d..b5496a2 100644
--- a/src/tint/lang/core/intrinsic/table_data.h
+++ b/src/tint/lang/core/intrinsic/table_data.h
@@ -634,7 +634,7 @@
         },
         /* print */
         [](MatchState* state, StyledText& out) {
-            out << style::Type << state->data[state->overload.templates + INDEX].name;
+            out << style::Type(state->data[state->overload.templates + INDEX].name);
         },
     };
 };
@@ -655,7 +655,7 @@
         },
         /* print */
         [](MatchState* state, StyledText& out) {
-            out << style::Variable << state->data[state->overload.templates + INDEX].name;
+            out << style::Variable(state->data[state->overload.templates + INDEX].name);
         },
     };
 };
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 3fd6ae4..3b96695 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -308,7 +308,7 @@
 
     for (auto& func : mod_.functions) {
         if (!all_functions_.Add(func.Get())) {
-            AddError(Source{}) << "function " << style::Function << Name(func.Get()) << style::Plain
+            AddError(Source{}) << "function " << style::Function(Name(func.Get()))
                                << " added to module multiple times";
         }
     }
diff --git a/src/tint/lang/core/texel_format.cc b/src/tint/lang/core/texel_format.cc
index c0f552a..1371c35 100644
--- a/src/tint/lang/core/texel_format.cc
+++ b/src/tint/lang/core/texel_format.cc
@@ -54,6 +54,9 @@
     if (str == "r32uint") {
         return TexelFormat::kR32Uint;
     }
+    if (str == "r8unorm") {
+        return TexelFormat::kR8Unorm;
+    }
     if (str == "rg32float") {
         return TexelFormat::kRg32Float;
     }
@@ -108,6 +111,8 @@
             return "r32sint";
         case TexelFormat::kR32Uint:
             return "r32uint";
+        case TexelFormat::kR8Unorm:
+            return "r8unorm";
         case TexelFormat::kRg32Float:
             return "rg32float";
         case TexelFormat::kRg32Sint:
diff --git a/src/tint/lang/core/texel_format.def b/src/tint/lang/core/texel_format.def
index 89e5f65..3d194a4 100644
--- a/src/tint/lang/core/texel_format.def
+++ b/src/tint/lang/core/texel_format.def
@@ -17,4 +17,7 @@
   rgba32uint
   rgba32sint
   rgba32float
+
+  // requires chromium_internal_graphite
+  r8unorm
 }
diff --git a/src/tint/lang/core/texel_format.h b/src/tint/lang/core/texel_format.h
index da6791d..6a9a4e2 100644
--- a/src/tint/lang/core/texel_format.h
+++ b/src/tint/lang/core/texel_format.h
@@ -50,6 +50,7 @@
     kR32Float,
     kR32Sint,
     kR32Uint,
+    kR8Unorm,
     kRg32Float,
     kRg32Sint,
     kRg32Uint,
@@ -83,9 +84,9 @@
 TexelFormat ParseTexelFormat(std::string_view str);
 
 constexpr std::string_view kTexelFormatStrings[] = {
-    "bgra8unorm", "r32float",    "r32sint",    "r32uint",    "rg32float",   "rg32sint",
-    "rg32uint",   "rgba16float", "rgba16sint", "rgba16uint", "rgba32float", "rgba32sint",
-    "rgba32uint", "rgba8sint",   "rgba8snorm", "rgba8uint",  "rgba8unorm",
+    "bgra8unorm", "r32float",   "r32sint",     "r32uint",    "r8unorm",    "rg32float",
+    "rg32sint",   "rg32uint",   "rgba16float", "rgba16sint", "rgba16uint", "rgba32float",
+    "rgba32sint", "rgba32uint", "rgba8sint",   "rgba8snorm", "rgba8uint",  "rgba8unorm",
 };
 
 }  // namespace tint::core
diff --git a/src/tint/lang/core/texel_format_bench.cc b/src/tint/lang/core/texel_format_bench.cc
index 90e7795..33f4455 100644
--- a/src/tint/lang/core/texel_format_bench.cc
+++ b/src/tint/lang/core/texel_format_bench.cc
@@ -45,30 +45,32 @@
 
 void TexelFormatParser(::benchmark::State& state) {
     const char* kStrings[] = {
-        "bgraunccrm",     "blranr3",      "bVra8unorm",   "bgra8unorm",   "bgra1unorm",
-        "bgrJqqnorm",     "bgr7ll8unorm", "qq32lppHat",   "c2fov",        "r32Goat",
-        "r32float",       "r3viiloat",    "r3WWflo8t",    "rxxfMoat",     "rXsingg",
-        "3siXt",          "r32s3nt",      "r32sint",      "E32sint",      "rPTTsint",
-        "r32sidxx",       "r442uint",     "r3SSuiVVt",    "R32R22t",      "r32uint",
-        "rFui9t",         "r32int",       "VOORRHnt",     "rgy2foat",     "l77nnrr2floGt",
-        "rg42fl00at",     "rg32float",    "rgoofat",      "rgzzflot",     "g11p2fliia",
-        "XXg32sint",      "rII39955nnnt", "aagHH2sinYSS", "rg32sint",     "rkk3it",
-        "gj3sRRn",        "r3bsnt",       "rg32jint",     "rg32unt",      "rgqint",
-        "rg32uint",       "rg32inNN",     "g3vvint",      "rg2uQQnt",     "rga16floft",
-        "rgja16float",    "rgNNww16f2oa", "rgba16float",  "rgba16flot",   "rgba16rrloat",
-        "rgGa16float",    "rgba16sFFnt",  "g16sEnt",      "rgb16rrint",   "rgba16sint",
-        "gba16sit",       "rXa1DsiJJt",   "rgasint",      "rg111kin",     "rgb16uint",
-        "rgJa16uit",      "rgba16uint",   "rgca16uint",   "rgba16Oint",   "KKgba__v6uintt",
-        "rgb832fxx5t",    "rgbaqq__lat",  "rgba32qloat",  "rgba32float",  "33gbO2floa66",
-        "rgboott6QQloat", "66ba32float",  "xba32zzinO6",  "ryyba32sint",  "rbZ32HinZ",
-        "rgba32sint",     "rgba3s4WWnq",  "rgba32sOOt",   "oogba2Yin",    "ga32unt",
-        "Fga32uint",      "rgb32uinw",    "rgba32uint",   "rgGf32uit",    "rgbaqKKuint",
-        "rFba32ummnt",    "rgba8snt",     "rgq8sint",     "rbbba8bin",    "rgba8sint",
-        "rgbisint",       "rgq8sinOO",    "rgbaTTvvint",  "rgFFa8snorm",  "rg00QsnPrm",
-        "rgbaPsnorm",     "rgba8snorm",   "rgb77ssnorm",  "rgba8snbbRRC", "rgbXX8snorm",
-        "qgCCOO8iOOt",    "rsauuinL",     "rgXa8uint",    "rgba8uint",    "rgba8int",
-        "rgbunqq",        "rg22a8uint",   "rybXX0nzzrm",  "rgVVa8iorP",   "rbaCunnnrm",
-        "rgba8unorm",     "ba8unoqqHHA",  "rga8unorm",    "rgfa8uKKo",
+        "bgraunccrm",    "blranr3",       "bVra8unorm",    "bgra8unorm",   "bgra1unorm",
+        "bgrJqqnorm",    "bgr7ll8unorm",  "qq32lppHat",    "c2fov",        "r32Goat",
+        "r32float",      "r3viiloat",     "r3WWflo8t",     "rxxfMoat",     "rXsingg",
+        "3siXt",         "r32s3nt",       "r32sint",       "E32sint",      "rPTTsint",
+        "r32sidxx",      "r442uint",      "r3SSuiVVt",     "R32R22t",      "r32uint",
+        "rFui9t",        "r32int",        "VOORRHnt",      "r8noym",       "lln8rrn77rm",
+        "r8un4r00",      "r8unorm",       "8noom",         "zznorm",       "riiuppo1",
+        "rg32floaXX",    "rg355flo99nII", "rrr32floSSaHH", "rg32float",    "kk3Hoat",
+        "jg2loaRR",      "rg2loab",       "rg32jint",      "rg32snt",      "rgqint",
+        "rg32sint",      "rg32inNN",      "g3vvint",       "rg2sQQnt",     "rg3furt",
+        "rgj2uint",      "rg3wNN82t",     "rg32uint",      "rg32unt",      "rgrr2uint",
+        "rG32uint",      "FFgba16float",  "rEb16loa",      "rga16frroat",  "rgba16float",
+        "rgba6loat",     "Xba1JJfDoat",   "rga6fl8a",      "rg111kin",     "rgb16sint",
+        "rgJa16sit",     "rgba16sint",    "rgca16sint",    "rgba16Oint",   "KKgba__v6sintt",
+        "rg5a16xxnt",    "__ba1uqqFt",    "rgbqq6uint",    "rgba16uint",   "33ba16u66nt",
+        "rtt6a1QQooint", "r66ba1uint",    "rgbx266loaOz",  "yygba32float", "gba3ZZHoat",
+        "rgba32float",   "WWga32floq4t",  "rgba3OOfoat",   "rYoha2flat",   "ga32snt",
+        "Fga32sint",     "rgb32sinw",     "rgba32sint",    "rgGf32sit",    "rgbaqKKsint",
+        "rFba32smmnt",   "rgba32uit",     "rqba3uint",     "rgbabb2uin",   "rgba32uint",
+        "rba32iint",     "qgba32uiOt",    "rgba32uiTTvv",  "rgbaFFsint",   "rgQa00siP",
+        "rgPa8sint",     "rgba8sint",     "rgssa77snt",    "Cgbbb8siRRt",  "rgba8sinXX",
+        "CqgbaOOsnorm",  "rgbu8ssrL",     "rgba8Xnorm",    "rgba8snorm",   "rgba8snrm",
+        "ba8sqqor",      "rgba8snor22",   "rzzyaXui0t",    "rPbi8uint",    "rgaCnnint",
+        "rgba8uint",     "bqqAAuinHH",    "rgba8unt",      "rgb8uiKf",     "raggunorm",
+        "rgb8unorm",     "rgba8uTNo4m",   "rgba8unorm",    "rgla8unppr7",  "rg8zznNNrm",
+        "bXXbauuuorm",
     };
     for (auto _ : state) {
         for (auto* str : kStrings) {
diff --git a/src/tint/lang/core/texel_format_test.cc b/src/tint/lang/core/texel_format_test.cc
index 1500656..ae49396 100644
--- a/src/tint/lang/core/texel_format_test.cc
+++ b/src/tint/lang/core/texel_format_test.cc
@@ -59,42 +59,43 @@
 static constexpr Case kValidCases[] = {
     {"bgra8unorm", TexelFormat::kBgra8Unorm},   {"r32float", TexelFormat::kR32Float},
     {"r32sint", TexelFormat::kR32Sint},         {"r32uint", TexelFormat::kR32Uint},
-    {"rg32float", TexelFormat::kRg32Float},     {"rg32sint", TexelFormat::kRg32Sint},
-    {"rg32uint", TexelFormat::kRg32Uint},       {"rgba16float", TexelFormat::kRgba16Float},
-    {"rgba16sint", TexelFormat::kRgba16Sint},   {"rgba16uint", TexelFormat::kRgba16Uint},
-    {"rgba32float", TexelFormat::kRgba32Float}, {"rgba32sint", TexelFormat::kRgba32Sint},
-    {"rgba32uint", TexelFormat::kRgba32Uint},   {"rgba8sint", TexelFormat::kRgba8Sint},
-    {"rgba8snorm", TexelFormat::kRgba8Snorm},   {"rgba8uint", TexelFormat::kRgba8Uint},
-    {"rgba8unorm", TexelFormat::kRgba8Unorm},
+    {"r8unorm", TexelFormat::kR8Unorm},         {"rg32float", TexelFormat::kRg32Float},
+    {"rg32sint", TexelFormat::kRg32Sint},       {"rg32uint", TexelFormat::kRg32Uint},
+    {"rgba16float", TexelFormat::kRgba16Float}, {"rgba16sint", TexelFormat::kRgba16Sint},
+    {"rgba16uint", TexelFormat::kRgba16Uint},   {"rgba32float", TexelFormat::kRgba32Float},
+    {"rgba32sint", TexelFormat::kRgba32Sint},   {"rgba32uint", TexelFormat::kRgba32Uint},
+    {"rgba8sint", TexelFormat::kRgba8Sint},     {"rgba8snorm", TexelFormat::kRgba8Snorm},
+    {"rgba8uint", TexelFormat::kRgba8Uint},     {"rgba8unorm", TexelFormat::kRgba8Unorm},
 };
 
 static constexpr Case kInvalidCases[] = {
-    {"bgraunccrm", TexelFormat::kUndefined},      {"blranr3", TexelFormat::kUndefined},
-    {"bVra8unorm", TexelFormat::kUndefined},      {"132float", TexelFormat::kUndefined},
-    {"32Jlqqat", TexelFormat::kUndefined},        {"ll3277loat", TexelFormat::kUndefined},
-    {"ppqq2snHH", TexelFormat::kUndefined},       {"r3cv", TexelFormat::kUndefined},
-    {"b2siGt", TexelFormat::kUndefined},          {"r32uiivt", TexelFormat::kUndefined},
-    {"8WW2uint", TexelFormat::kUndefined},        {"rxxuint", TexelFormat::kUndefined},
-    {"rX2flggat", TexelFormat::kUndefined},       {"rg3XVut", TexelFormat::kUndefined},
-    {"3g32float", TexelFormat::kUndefined},       {"rg3Esint", TexelFormat::kUndefined},
-    {"PP32TTint", TexelFormat::kUndefined},       {"xxg32ddnt", TexelFormat::kUndefined},
-    {"44g32uint", TexelFormat::kUndefined},       {"rSS32uinVV", TexelFormat::kUndefined},
-    {"R322Rint", TexelFormat::kUndefined},        {"rgba16fF9a", TexelFormat::kUndefined},
-    {"rgba16floa", TexelFormat::kUndefined},      {"rOObVR16floH", TexelFormat::kUndefined},
-    {"ryba1sint", TexelFormat::kUndefined},       {"r77ba1nnsllrrt", TexelFormat::kUndefined},
-    {"rgb4006sint", TexelFormat::kUndefined},     {"rb1uioot", TexelFormat::kUndefined},
-    {"rga1uzznt", TexelFormat::kUndefined},       {"r11b1uppiit", TexelFormat::kUndefined},
-    {"rgba32fXXoat", TexelFormat::kUndefined},    {"rgbII99355float", TexelFormat::kUndefined},
-    {"rgbaa32fSSrHHYt", TexelFormat::kUndefined}, {"rbkk2Hit", TexelFormat::kUndefined},
-    {"jgba3sgRR", TexelFormat::kUndefined},       {"rgbab2si", TexelFormat::kUndefined},
-    {"rgba32jint", TexelFormat::kUndefined},      {"rba32uint", TexelFormat::kUndefined},
-    {"rgba2uqn", TexelFormat::kUndefined},        {"rbaNNsint", TexelFormat::kUndefined},
-    {"rga8invv", TexelFormat::kUndefined},        {"gba8sQQnt", TexelFormat::kUndefined},
-    {"rgbsnrrff", TexelFormat::kUndefined},       {"rgba8snojm", TexelFormat::kUndefined},
-    {"NNgba8sww2m", TexelFormat::kUndefined},     {"rgba8uit", TexelFormat::kUndefined},
-    {"rrgba8uint", TexelFormat::kUndefined},      {"rgba8uiGt", TexelFormat::kUndefined},
-    {"rgba8unFFrm", TexelFormat::kUndefined},     {"g8unErm", TexelFormat::kUndefined},
-    {"rgb8urrorm", TexelFormat::kUndefined},
+    {"bgraunccrm", TexelFormat::kUndefined},     {"blranr3", TexelFormat::kUndefined},
+    {"bVra8unorm", TexelFormat::kUndefined},     {"132float", TexelFormat::kUndefined},
+    {"32Jlqqat", TexelFormat::kUndefined},       {"ll3277loat", TexelFormat::kUndefined},
+    {"ppqq2snHH", TexelFormat::kUndefined},      {"r3cv", TexelFormat::kUndefined},
+    {"b2siGt", TexelFormat::kUndefined},         {"r32uiivt", TexelFormat::kUndefined},
+    {"8WW2uint", TexelFormat::kUndefined},       {"rxxuint", TexelFormat::kUndefined},
+    {"rXnorgg", TexelFormat::kUndefined},        {"8noXm", TexelFormat::kUndefined},
+    {"r8un3rm", TexelFormat::kUndefined},        {"rg32floEt", TexelFormat::kUndefined},
+    {"rgTTP2loat", TexelFormat::kUndefined},     {"ddg32loxxt", TexelFormat::kUndefined},
+    {"44g32sint", TexelFormat::kUndefined},      {"rSS32sinVV", TexelFormat::kUndefined},
+    {"R322Rint", TexelFormat::kUndefined},       {"gF29int", TexelFormat::kUndefined},
+    {"rg3uint", TexelFormat::kUndefined},        {"ROOgH2iVt", TexelFormat::kUndefined},
+    {"rgya16flot", TexelFormat::kUndefined},     {"rrgba16flll77Gt", TexelFormat::kUndefined},
+    {"rgba46fl00at", TexelFormat::kUndefined},   {"rb1sioot", TexelFormat::kUndefined},
+    {"rga1szznt", TexelFormat::kUndefined},      {"r11b1sppiit", TexelFormat::kUndefined},
+    {"XXgba16uint", TexelFormat::kUndefined},    {"IIgb9916nni55t", TexelFormat::kUndefined},
+    {"rYbaSSrruiHHat", TexelFormat::kUndefined}, {"gHkkfloat", TexelFormat::kUndefined},
+    {"jRRba3fgat", TexelFormat::kUndefined},     {"rba32flbt", TexelFormat::kUndefined},
+    {"rgba32jint", TexelFormat::kUndefined},     {"rba32sint", TexelFormat::kUndefined},
+    {"rgba2sqn", TexelFormat::kUndefined},       {"rgba3NNuit", TexelFormat::kUndefined},
+    {"rga3vvint", TexelFormat::kUndefined},      {"rgba32uinQ", TexelFormat::kUndefined},
+    {"rgbasirf", TexelFormat::kUndefined},       {"rgbajsint", TexelFormat::kUndefined},
+    {"wNNgbasin2", TexelFormat::kUndefined},     {"rgba8snrm", TexelFormat::kUndefined},
+    {"rgba8srrorm", TexelFormat::kUndefined},    {"rgba8Gnorm", TexelFormat::kUndefined},
+    {"rgba8uFFnt", TexelFormat::kUndefined},     {"Ega8un", TexelFormat::kUndefined},
+    {"rgrr8uint", TexelFormat::kUndefined},      {"gba8unom", TexelFormat::kUndefined},
+    {"rXa8DnoJJm", TexelFormat::kUndefined},     {"rganorm", TexelFormat::kUndefined},
 };
 
 using TexelFormatParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/core/type/storage_texture.cc b/src/tint/lang/core/type/storage_texture.cc
index 879ff99..ea7c9ca 100644
--- a/src/tint/lang/core/type/storage_texture.cc
+++ b/src/tint/lang/core/type/storage_texture.cc
@@ -80,6 +80,7 @@
             return type_mgr.Get<I32>();
         }
 
+        case core::TexelFormat::kR8Unorm:
         case core::TexelFormat::kBgra8Unorm:
         case core::TexelFormat::kRgba8Unorm:
         case core::TexelFormat::kRgba8Snorm:
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 6ecc839..21ff1b8 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -274,8 +274,9 @@
             "GLSL", builder_.AST(), diagnostics_,
             Vector{
                 wgsl::Extension::kChromiumDisableUniformityAnalysis,
-                wgsl::Extension::kChromiumInternalDualSourceBlending,
                 wgsl::Extension::kChromiumExperimentalPushConstant,
+                wgsl::Extension::kChromiumInternalDualSourceBlending,
+                wgsl::Extension::kChromiumInternalGraphite,
                 wgsl::Extension::kF16,
             })) {
         return false;
@@ -2028,6 +2029,9 @@
             case core::TexelFormat::kRgba32Float:
                 out << "rgba32f";
                 break;
+            case core::TexelFormat::kR8Unorm:
+                out << "r8";
+                break;
             case core::TexelFormat::kUndefined:
                 TINT_ICE() << "invalid texel format";
                 return;
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 658beeb..387ddb9 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -105,8 +105,9 @@
 
 const char kTempNamePrefix[] = "tint_tmp";
 
-const char* image_format_to_rwtexture_type(core::TexelFormat image_format) {
+const char* ImageFormatToRWtextureType(core::TexelFormat image_format) {
     switch (image_format) {
+        case core::TexelFormat::kR8Unorm:
         case core::TexelFormat::kBgra8Unorm:
         case core::TexelFormat::kRgba8Unorm:
         case core::TexelFormat::kRgba8Snorm:
@@ -388,11 +389,12 @@
             "HLSL", builder_.AST(), diagnostics_,
             Vector{
                 wgsl::Extension::kChromiumDisableUniformityAnalysis,
+                wgsl::Extension::kChromiumExperimentalPixelLocal,
                 wgsl::Extension::kChromiumExperimentalPushConstant,
                 wgsl::Extension::kChromiumExperimentalSubgroups,
-                wgsl::Extension::kF16,
                 wgsl::Extension::kChromiumInternalDualSourceBlending,
-                wgsl::Extension::kChromiumExperimentalPixelLocal,
+                wgsl::Extension::kChromiumInternalGraphite,
+                wgsl::Extension::kF16,
             })) {
         return false;
     }
@@ -3459,7 +3461,7 @@
             return false;
         }
         out << "RasterizerOrderedTexture2D";
-        auto* component = image_format_to_rwtexture_type(storage->texel_format());
+        auto* component = ImageFormatToRWtextureType(storage->texel_format());
         if (TINT_UNLIKELY(!component)) {
             TINT_ICE() << "Unsupported StorageTexture TexelFormat: "
                        << static_cast<int>(storage->texel_format());
@@ -4460,7 +4462,7 @@
             }
 
             if (storage) {
-                auto* component = image_format_to_rwtexture_type(storage->texel_format());
+                auto* component = ImageFormatToRWtextureType(storage->texel_format());
                 if (TINT_UNLIKELY(!component)) {
                     TINT_ICE() << "Unsupported StorageTexture TexelFormat: "
                                << static_cast<int>(storage->texel_format());
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index c4f7487..3baa716 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -83,7 +83,7 @@
     return BuildU32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "u32";
+    out << style::Type("u32");
   }
 };
 
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index c2442f5..9a574b7 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -273,10 +273,11 @@
             "MSL", builder_.AST(), diagnostics_,
             Vector{
                 wgsl::Extension::kChromiumDisableUniformityAnalysis,
+                wgsl::Extension::kChromiumExperimentalFramebufferFetch,
                 wgsl::Extension::kChromiumExperimentalPixelLocal,
                 wgsl::Extension::kChromiumExperimentalSubgroups,
-                wgsl::Extension::kChromiumExperimentalFramebufferFetch,
                 wgsl::Extension::kChromiumInternalDualSourceBlending,
+                wgsl::Extension::kChromiumInternalGraphite,
                 wgsl::Extension::kChromiumInternalRelaxedUniformLayout,
                 wgsl::Extension::kF16,
             })) {
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index a78f3ee..5a1e8c0 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -178,7 +178,7 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x2* const tint_symbol) {
-  {
+  if ((local_idx < 1u)) {
     *(tint_symbol) = float2x2(float2(0.0f), float2(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -300,7 +300,7 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup S2* const tint_symbol_1) {
-  {
+  if ((local_idx < 1u)) {
     S2 const tint_symbol = S2{};
     *(tint_symbol_1) = tint_symbol;
   }
@@ -397,7 +397,7 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x2* const tint_symbol, threadgroup float2x3* const tint_symbol_1, threadgroup float2x4* const tint_symbol_2) {
-  {
+  if ((local_idx < 1u)) {
     *(tint_symbol) = float2x2(float2(0.0f), float2(0.0f));
     *(tint_symbol_1) = float2x3(float3(0.0f), float3(0.0f));
     *(tint_symbol_2) = float2x4(float4(0.0f), float4(0.0f));
@@ -406,7 +406,7 @@
 }
 
 void tint_zero_workgroup_memory_1(uint local_idx_1, threadgroup float3x2* const tint_symbol_3, threadgroup float3x3* const tint_symbol_4, threadgroup float3x4* const tint_symbol_5) {
-  {
+  if ((local_idx_1 < 1u)) {
     *(tint_symbol_3) = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
     *(tint_symbol_4) = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
     *(tint_symbol_5) = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
@@ -415,7 +415,7 @@
 }
 
 void tint_zero_workgroup_memory_2(uint local_idx_2, threadgroup float4x2* const tint_symbol_6, threadgroup float4x3* const tint_symbol_7, threadgroup float4x4* const tint_symbol_8) {
-  {
+  if ((local_idx_2 < 1u)) {
     *(tint_symbol_6) = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
     *(tint_symbol_7) = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
     *(tint_symbol_8) = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 691263a..b00c77c 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -84,7 +84,7 @@
     return BuildBool(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "bool";
+    out << style::Type("bool");
   }
 };
 
@@ -98,7 +98,7 @@
     return BuildF32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f32";
+    out << style::Type("f32");
   }
 };
 
@@ -112,7 +112,7 @@
     return BuildF16(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f16";
+    out << style::Type("f16");
   }
 };
 
@@ -126,7 +126,7 @@
     return BuildI32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "i32";
+    out << style::Type("i32");
   }
 };
 
@@ -140,7 +140,7 @@
     return BuildU32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "u32";
+    out << style::Type("u32");
   }
 };
 
@@ -160,7 +160,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec2", "<", T, ">");
   }
 };
 
@@ -180,7 +180,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec3", "<", T, ">");
   }
 };
 
@@ -200,7 +200,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec4", "<", T, ">");
   }
 };
 
@@ -220,7 +220,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x2", "<", T, ">");
   }
 };
 
@@ -240,7 +240,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x3", "<", T, ">");
   }
 };
 
@@ -260,7 +260,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x4", "<", T, ">");
   }
 };
 
@@ -280,7 +280,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x2", "<", T, ">");
   }
 };
 
@@ -300,7 +300,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x3", "<", T, ">");
   }
 };
 
@@ -320,7 +320,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x4", "<", T, ">");
   }
 };
 
@@ -340,7 +340,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x2", "<", T, ">");
   }
 };
 
@@ -360,7 +360,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x3", "<", T, ">");
   }
 };
 
@@ -380,7 +380,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x4", "<", T, ">");
   }
 };
 
@@ -406,7 +406,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "vec" << style::Type << N << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("vec", N, "<", T, ">");
   }
 };
 
@@ -438,7 +438,7 @@
   state->PrintNum(N);StyledText M;
   state->PrintNum(M);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "mat" << style::Type << N << style::Type << "x" << style::Type << M << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("mat", N, "x", M, "<", T, ">");
   }
 };
 
@@ -458,7 +458,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "atomic" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("atomic", "<", T, ">");
   }
 };
 
@@ -472,7 +472,7 @@
     return BuildSampler(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler";
+    out << style::Type("sampler");
   }
 };
 
@@ -486,7 +486,7 @@
     return BuildSamplerComparison(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler_comparison";
+    out << style::Type("sampler_comparison");
   }
 };
 
@@ -506,7 +506,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_1d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_1d", "<", T, ">");
   }
 };
 
@@ -526,7 +526,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d", "<", T, ">");
   }
 };
 
@@ -546,7 +546,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d_array", "<", T, ">");
   }
 };
 
@@ -566,7 +566,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_3d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_3d", "<", T, ">");
   }
 };
 
@@ -586,7 +586,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube", "<", T, ">");
   }
 };
 
@@ -606,7 +606,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube_array", "<", T, ">");
   }
 };
 
@@ -626,7 +626,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_multisampled_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_multisampled_2d", "<", T, ">");
   }
 };
 
@@ -640,7 +640,7 @@
     return BuildTextureDepth2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d";
+    out << style::Type("texture_depth_2d");
   }
 };
 
@@ -654,7 +654,7 @@
     return BuildTextureDepth2DArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d_array";
+    out << style::Type("texture_depth_2d_array");
   }
 };
 
@@ -668,7 +668,7 @@
     return BuildTextureDepthCube(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube";
+    out << style::Type("texture_depth_cube");
   }
 };
 
@@ -682,7 +682,7 @@
     return BuildTextureDepthCubeArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube_array";
+    out << style::Type("texture_depth_cube_array");
   }
 };
 
@@ -696,7 +696,7 @@
     return BuildTextureDepthMultisampled2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_multisampled_2d";
+    out << style::Type("texture_depth_multisampled_2d");
   }
 };
 
@@ -722,7 +722,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_1d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_1d", "<", F, ", ", A, ">");
   }
 };
 
@@ -748,7 +748,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d", "<", F, ", ", A, ">");
   }
 };
 
@@ -774,7 +774,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d_array" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d_array", "<", F, ", ", A, ">");
   }
 };
 
@@ -800,7 +800,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_3d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_3d", "<", F, ", ", A, ">");
   }
 };
 
@@ -832,7 +832,7 @@
   state->PrintNum(S);StyledText T;
   state->PrintType(T);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "ptr" << style::Code << "<" << style::Type << S << style::Code << ", " << style::Type << T << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("ptr", "<", S, ", ", T, ", ", A, ">");
   }
 };
 
@@ -846,7 +846,7 @@
     return BuildStructWithRuntimeArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "struct_with_runtime_array";
+    out << style::Type("struct_with_runtime_array");
   }
 };
 
@@ -866,7 +866,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "sampled_image" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("sampled_image", "<", T, ">");
   }
 };
 
@@ -885,7 +885,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match iu32'
@@ -902,7 +902,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fiu32'
@@ -922,7 +922,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar'
@@ -948,7 +948,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match samplers'
@@ -965,7 +965,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kSamplerMatcher.print(nullptr, out); out << TextStyle{} << " or "; kSamplerComparisonMatcher.print(nullptr, out);}
+ kSamplerMatcher.print(nullptr, out); out << style::Plain(" or "); kSamplerComparisonMatcher.print(nullptr, out);}
 };
 
 /// EnumMatcher for 'match read_write'
@@ -977,7 +977,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read_write";
+  out<< style::Enum("read_write");
   }
 };
 
@@ -990,7 +990,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "storage";
+  out<< style::Enum("storage");
   }
 };
 
@@ -1006,7 +1006,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "workgroup"<< TextStyle{} << " or " << style::Enum << "storage";
+  out<< style::Enum("workgroup")<< style::Plain(" or ") << style::Enum("storage");
   }
 };
 
@@ -1014,6 +1014,7 @@
 constexpr NumberMatcher kF32TexelFormatMatcher {
 /* match */ [](MatchState&, Number number) -> Number {
     switch (static_cast<core::TexelFormat>(number.Value())) {
+      case core::TexelFormat::kR8Unorm:
       case core::TexelFormat::kBgra8Unorm:
       case core::TexelFormat::kRgba8Unorm:
       case core::TexelFormat::kRgba8Snorm:
@@ -1027,7 +1028,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "bgra8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8snorm"<< TextStyle{} << ", " << style::Enum << "rgba16float"<< TextStyle{} << ", " << style::Enum << "r32float"<< TextStyle{} << ", " << style::Enum << "rg32float"<< TextStyle{} << " or " << style::Enum << "rgba32float";
+  out<< style::Enum("r8unorm")<< style::Plain(", ") << style::Enum("bgra8unorm")<< style::Plain(", ") << style::Enum("rgba8unorm")<< style::Plain(", ") << style::Enum("rgba8snorm")<< style::Plain(", ") << style::Enum("rgba16float")<< style::Plain(", ") << style::Enum("r32float")<< style::Plain(", ") << style::Enum("rg32float")<< style::Plain(" or ") << style::Enum("rgba32float");
   }
 };
 
@@ -1046,7 +1047,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8sint"<< TextStyle{} << ", " << style::Enum << "rgba16sint"<< TextStyle{} << ", " << style::Enum << "r32sint"<< TextStyle{} << ", " << style::Enum << "rg32sint"<< TextStyle{} << " or " << style::Enum << "rgba32sint";
+  out<< style::Enum("rgba8sint")<< style::Plain(", ") << style::Enum("rgba16sint")<< style::Plain(", ") << style::Enum("r32sint")<< style::Plain(", ") << style::Enum("rg32sint")<< style::Plain(" or ") << style::Enum("rgba32sint");
   }
 };
 
@@ -1065,7 +1066,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8uint"<< TextStyle{} << ", " << style::Enum << "rgba16uint"<< TextStyle{} << ", " << style::Enum << "r32uint"<< TextStyle{} << ", " << style::Enum << "rg32uint"<< TextStyle{} << " or " << style::Enum << "rgba32uint";
+  out<< style::Enum("rgba8uint")<< style::Plain(", ") << style::Enum("rgba16uint")<< style::Plain(", ") << style::Enum("r32uint")<< style::Plain(", ") << style::Enum("rg32uint")<< style::Plain(" or ") << style::Enum("rgba32uint");
   }
 };
 
@@ -1081,7 +1082,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("read")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
@@ -1097,7 +1098,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "write"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("write")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 96fb6c4..9dabb13 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -99,7 +99,8 @@
   | sampler_comparison
 
 match f32_texel_format
-  : texel_format.bgra8unorm
+  : texel_format.r8unorm
+  | texel_format.bgra8unorm
   | texel_format.rgba8unorm
   | texel_format.rgba8snorm
   | texel_format.rgba16float
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index a412f49..9673668 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -279,8 +279,9 @@
                 wgsl::Extension::kChromiumDisableUniformityAnalysis,
                 wgsl::Extension::kChromiumExperimentalPushConstant,
                 wgsl::Extension::kChromiumExperimentalSubgroups,
-                wgsl::Extension::kF16,
                 wgsl::Extension::kChromiumInternalDualSourceBlending,
+                wgsl::Extension::kChromiumInternalGraphite,
+                wgsl::Extension::kF16,
             })) {
         return false;
     }
@@ -4109,6 +4110,9 @@
         case core::TexelFormat::kBgra8Unorm:
             TINT_ICE() << "bgra8unorm should have been polyfilled to rgba8unorm";
             return SpvImageFormatUnknown;
+        case core::TexelFormat::kR8Unorm:
+            module_.PushCapability(SpvCapabilityStorageImageExtendedFormats);
+            return SpvImageFormatR8;
         case core::TexelFormat::kR32Uint:
             return SpvImageFormatR32ui;
         case core::TexelFormat::kR32Sint:
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 9fc98fd..c38f018 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -2228,6 +2228,9 @@
             case core::TexelFormat::kBgra8Unorm:
                 TINT_ICE() << "bgra8unorm should have been polyfilled to rgba8unorm";
                 return SpvImageFormatUnknown;
+            case core::TexelFormat::kR8Unorm:
+                module_.PushCapability(SpvCapabilityStorageImageExtendedFormats);
+                return SpvImageFormatR8;
             case core::TexelFormat::kR32Uint:
                 return SpvImageFormatR32ui;
             case core::TexelFormat::kR32Sint:
diff --git a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
index d0e7f74..7745eb3 100644
--- a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.cc
@@ -219,6 +219,10 @@
 
             // Determine the block type used to emit these statements.
 
+            // TODO(crbug.com/tint/2143): Always emit an if statement around zero init, even when
+            // workgroup size matches num_iteration, to work around bugs in certain drivers.
+            constexpr bool kWorkaroundUnconditionalZeroInitDriverBug = true;
+
             if (workgroup_size_const == 0 || num_iterations > workgroup_size_const) {
                 // Either the workgroup size is dynamic, or smaller than num_iterations.
                 // In either case, we need to generate a for loop to ensure we
@@ -243,14 +247,14 @@
                 }
                 auto* for_loop = b.For(init, cond, cont, b.Block(block));
                 init_body.Push(for_loop);
-            } else if (num_iterations < workgroup_size_const) {
+            } else if (num_iterations < workgroup_size_const ||
+                       kWorkaroundUnconditionalZeroInitDriverBug) {
                 // Workgroup size is a known constant, but is greater than
                 // num_iterations. Emit an if statement:
                 //
                 //  if (local_index < num_iterations) {
                 //    ...
                 //  }
-
                 auto* cond = b.LessThan(local_idx, u32(num_iterations));
                 auto block = DeclareArrayIndices(num_iterations, array_indices,
                                                  [&] { return b.Expr(local_idx); });
diff --git a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_test.cc b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_test.cc
index 440e2a4..c5b8677 100644
--- a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_test.cc
@@ -139,7 +139,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -170,7 +170,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -205,7 +205,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -245,7 +245,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -281,7 +281,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -312,7 +312,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -354,7 +354,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     a = i32();
     b.x = i32();
   }
@@ -421,7 +421,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     a = i32();
     b.x = i32();
   }
@@ -720,7 +720,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     a = i32();
     b.x = i32();
   }
@@ -787,7 +787,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     a = i32();
     b.x = i32();
   }
@@ -864,7 +864,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     a = i32();
   }
   for(var idx : u32 = local_idx; (idx < 32u); idx = (idx + 1u)) {
@@ -975,7 +975,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     a = i32();
   }
   for(var idx : u32 = local_idx; (idx < 32u); idx = (idx + 1u)) {
@@ -1073,7 +1073,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -1120,7 +1120,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx_1 : u32) {
-  {
+  if ((local_idx_1 < 1u)) {
     v = i32();
   }
   workgroupBarrier();
@@ -1161,7 +1161,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     atomicStore(&(i), i32());
     atomicStore(&(u), u32());
   }
@@ -1198,7 +1198,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     atomicStore(&(i), i32());
     atomicStore(&(u), u32());
   }
@@ -1241,7 +1241,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     w.a = i32();
     atomicStore(&(w.i), i32());
     w.b = f32();
@@ -1292,7 +1292,7 @@
 )";
     auto* expect = R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     w.a = i32();
     atomicStore(&(w.i), i32());
     w.b = f32();
@@ -1528,7 +1528,7 @@
     auto* expect =
         R"(
 fn tint_zero_workgroup_memory(local_idx : u32) {
-  {
+  if ((local_idx < 1u)) {
     W = mat2x2<f32>();
   }
   workgroupBarrier();
diff --git a/src/tint/lang/wgsl/helpers/check_supported_extensions.cc b/src/tint/lang/wgsl/helpers/check_supported_extensions.cc
index 72f65c6..5a4372c 100644
--- a/src/tint/lang/wgsl/helpers/check_supported_extensions.cc
+++ b/src/tint/lang/wgsl/helpers/check_supported_extensions.cc
@@ -49,8 +49,8 @@
         for (auto* ext : enable->extensions) {
             if (!set.Contains(ext->name)) {
                 diags.AddError(diag::System::Writer, ext->source)
-                    << writer_name << " backend does not support extension " << style::Code
-                    << ext->name;
+                    << writer_name << " backend does not support extension "
+                    << style::Code(ext->name);
                 return false;
             }
         }
diff --git a/src/tint/lang/wgsl/inspector/resource_binding.cc b/src/tint/lang/wgsl/inspector/resource_binding.cc
index 810e544..0cf21bb 100644
--- a/src/tint/lang/wgsl/inspector/resource_binding.cc
+++ b/src/tint/lang/wgsl/inspector/resource_binding.cc
@@ -120,6 +120,8 @@
             return ResourceBinding::TexelFormat::kRgba32Sint;
         case core::TexelFormat::kRgba32Float:
             return ResourceBinding::TexelFormat::kRgba32Float;
+        case core::TexelFormat::kR8Unorm:
+            return ResourceBinding::TexelFormat::kR8Unorm;
         case core::TexelFormat::kUndefined:
             return ResourceBinding::TexelFormat::kNone;
     }
diff --git a/src/tint/lang/wgsl/inspector/resource_binding.h b/src/tint/lang/wgsl/inspector/resource_binding.h
index 86fb720..1c1f801 100644
--- a/src/tint/lang/wgsl/inspector/resource_binding.h
+++ b/src/tint/lang/wgsl/inspector/resource_binding.h
@@ -80,6 +80,7 @@
         kRgba32Uint,
         kRgba32Sint,
         kRgba32Float,
+        kR8Unorm,
         kNone,
     };
 
diff --git a/src/tint/lang/wgsl/intrinsic/data.cc b/src/tint/lang/wgsl/intrinsic/data.cc
index 96b5ad0..3a84f91 100644
--- a/src/tint/lang/wgsl/intrinsic/data.cc
+++ b/src/tint/lang/wgsl/intrinsic/data.cc
@@ -106,7 +106,7 @@
     return BuildBool(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "bool";
+    out << style::Type("bool");
   }
 };
 
@@ -120,7 +120,7 @@
     return BuildIa(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out  << style::Type << "abstract-int";
+    out << style::Type("abstract-int");
   }
 };
 
@@ -134,7 +134,7 @@
     return BuildFa(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out  << style::Type << "abstract-float";
+    out << style::Type("abstract-float");
   }
 };
 
@@ -148,7 +148,7 @@
     return BuildI32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "i32";
+    out << style::Type("i32");
   }
 };
 
@@ -162,7 +162,7 @@
     return BuildU32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "u32";
+    out << style::Type("u32");
   }
 };
 
@@ -176,7 +176,7 @@
     return BuildF32(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f32";
+    out << style::Type("f32");
   }
 };
 
@@ -190,7 +190,7 @@
     return BuildF16(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "f16";
+    out << style::Type("f16");
   }
 };
 
@@ -210,7 +210,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec2", "<", T, ">");
   }
 };
 
@@ -230,7 +230,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec3", "<", T, ">");
   }
 };
 
@@ -250,7 +250,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "vec4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("vec4", "<", T, ">");
   }
 };
 
@@ -270,7 +270,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x2", "<", T, ">");
   }
 };
 
@@ -290,7 +290,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x3", "<", T, ">");
   }
 };
 
@@ -310,7 +310,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat2x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat2x4", "<", T, ">");
   }
 };
 
@@ -330,7 +330,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x2", "<", T, ">");
   }
 };
 
@@ -350,7 +350,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x3", "<", T, ">");
   }
 };
 
@@ -370,7 +370,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat3x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat3x4", "<", T, ">");
   }
 };
 
@@ -390,7 +390,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x2" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x2", "<", T, ">");
   }
 };
 
@@ -410,7 +410,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x3", "<", T, ">");
   }
 };
 
@@ -430,7 +430,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "mat4x4" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("mat4x4", "<", T, ">");
   }
 };
 
@@ -456,7 +456,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "vec" << style::Type << N << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("vec", N, "<", T, ">");
   }
 };
 
@@ -488,7 +488,7 @@
   state->PrintNum(N);StyledText M;
   state->PrintNum(M);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "mat" << style::Type << N << style::Type << "x" << style::Type << M << style::Type << "<" << style::Type << T << style::Type << ">";
+    out << style::Type("mat", N, "x", M, "<", T, ">");
   }
 };
 
@@ -520,7 +520,7 @@
   state->PrintNum(S);StyledText T;
   state->PrintType(T);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "ptr" << style::Code << "<" << style::Type << S << style::Code << ", " << style::Type << T << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("ptr", "<", S, ", ", T, ", ", A, ">");
   }
 };
 
@@ -540,7 +540,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "atomic" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("atomic", "<", T, ">");
   }
 };
 
@@ -560,7 +560,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("array", "<", T, ">");
   }
 };
 
@@ -574,7 +574,7 @@
     return BuildSampler(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler";
+    out << style::Type("sampler");
   }
 };
 
@@ -588,7 +588,7 @@
     return BuildSamplerComparison(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "sampler_comparison";
+    out << style::Type("sampler_comparison");
   }
 };
 
@@ -608,7 +608,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_1d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_1d", "<", T, ">");
   }
 };
 
@@ -628,7 +628,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d", "<", T, ">");
   }
 };
 
@@ -648,7 +648,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_2d_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_2d_array", "<", T, ">");
   }
 };
 
@@ -668,7 +668,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_3d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_3d", "<", T, ">");
   }
 };
 
@@ -688,7 +688,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube", "<", T, ">");
   }
 };
 
@@ -708,7 +708,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_cube_array" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_cube_array", "<", T, ">");
   }
 };
 
@@ -728,7 +728,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "texture_multisampled_2d" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("texture_multisampled_2d", "<", T, ">");
   }
 };
 
@@ -742,7 +742,7 @@
     return BuildTextureDepth2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d";
+    out << style::Type("texture_depth_2d");
   }
 };
 
@@ -756,7 +756,7 @@
     return BuildTextureDepth2DArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_2d_array";
+    out << style::Type("texture_depth_2d_array");
   }
 };
 
@@ -770,7 +770,7 @@
     return BuildTextureDepthCube(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube";
+    out << style::Type("texture_depth_cube");
   }
 };
 
@@ -784,7 +784,7 @@
     return BuildTextureDepthCubeArray(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_cube_array";
+    out << style::Type("texture_depth_cube_array");
   }
 };
 
@@ -798,7 +798,7 @@
     return BuildTextureDepthMultisampled2D(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_depth_multisampled_2d";
+    out << style::Type("texture_depth_multisampled_2d");
   }
 };
 
@@ -824,7 +824,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_1d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_1d", "<", F, ", ", A, ">");
   }
 };
 
@@ -850,7 +850,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d", "<", F, ", ", A, ">");
   }
 };
 
@@ -876,7 +876,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_2d_array" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_2d_array", "<", F, ", ", A, ">");
   }
 };
 
@@ -902,7 +902,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText F;
   state->PrintNum(F);StyledText A;
   state->PrintNum(A);
-    out << style::Type << "texture_storage_3d" << style::Code << "<" << style::Type << F << style::Code << ", " << style::Type << A << style::Code << ">";
+    out << style::Type("texture_storage_3d", "<", F, ", ", A, ">");
   }
 };
 
@@ -916,7 +916,7 @@
     return BuildTextureExternal(state, ty);
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
-    out << style::Type << "texture_external";
+    out << style::Type("texture_external");
   }
 };
 
@@ -936,7 +936,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "packedVec3" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("packedVec3", "<", T, ">");
   }
 };
 
@@ -956,7 +956,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__modf_result_" << style::Type << T;
+    out << style::Type("__modf_result_", T);
   }
 };
 
@@ -982,7 +982,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__modf_result_vec" << style::Type << N << style::Type << "_" << style::Type << T;
+    out << style::Type("__modf_result_vec", N, "_", T);
   }
 };
 
@@ -1002,7 +1002,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__frexp_result_" << style::Type << T;
+    out << style::Type("__frexp_result_", T);
   }
 };
 
@@ -1028,7 +1028,7 @@
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText N;
   state->PrintNum(N);StyledText T;
   state->PrintType(T);
-    out  << style::Type << "__frexp_result_vec" << style::Type << N << style::Type << "_" << style::Type << T;
+    out << style::Type("__frexp_result_vec", N, "_", T);
   }
 };
 
@@ -1048,7 +1048,7 @@
   },
 /* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
   state->PrintType(T);
-    out << style::Type << "__atomic_compare_exchange_result" << style::Code << "<" << style::Type << T << style::Code << ">";
+    out << style::Type("__atomic_compare_exchange_result", "<", T, ">");
   }
 };
 
@@ -1082,7 +1082,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match concrete_scalar'
@@ -1108,7 +1108,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_f32'
@@ -1137,7 +1137,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_f16'
@@ -1166,7 +1166,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_i32'
@@ -1195,7 +1195,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_u32'
@@ -1224,7 +1224,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kBoolMatcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kBoolMatcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match scalar_no_bool'
@@ -1253,7 +1253,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kF16Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kF16Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fia_fiu32_f16'
@@ -1282,7 +1282,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kFaMatcher.print(nullptr, out); out << style::Plain(", "); kIaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fia_fi32_f16'
@@ -1308,7 +1308,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kFaMatcher.print(nullptr, out); out << style::Plain(", "); kIaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fia_fiu32'
@@ -1334,7 +1334,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kFaMatcher.print(nullptr, out); out << style::Plain(", "); kIaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fa_f32'
@@ -1351,7 +1351,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kFaMatcher.print(nullptr, out); out << TextStyle{} << " or "; kF32Matcher.print(nullptr, out);}
+ kFaMatcher.print(nullptr, out); out << style::Plain(" or "); kF32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fa_f32_f16'
@@ -1371,7 +1371,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kFaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kFaMatcher.print(nullptr, out); out << style::Plain(", "); kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match ia_iu32'
@@ -1391,7 +1391,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match ia_i32'
@@ -1408,7 +1408,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kIaMatcher.print(nullptr, out); out << TextStyle{} << " or "; kI32Matcher.print(nullptr, out);}
+ kIaMatcher.print(nullptr, out); out << style::Plain(" or "); kI32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fiu32_f16'
@@ -1431,7 +1431,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kU32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(", "); kU32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fiu32'
@@ -1451,7 +1451,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fi32_f16'
@@ -1471,7 +1471,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << ", "; kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(", "); kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match fi32'
@@ -1488,7 +1488,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kI32Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kI32Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match f32_f16'
@@ -1505,7 +1505,7 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kF32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kF16Matcher.print(nullptr, out);}
+ kF32Matcher.print(nullptr, out); out << style::Plain(" or "); kF16Matcher.print(nullptr, out);}
 };
 
 /// TypeMatcher for 'match iu32'
@@ -1522,13 +1522,14 @@
 /* print */ [](MatchState*, StyledText& out) {
     // Note: We pass nullptr to the Matcher.print() functions, as matchers do not support
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
- kI32Matcher.print(nullptr, out); out << TextStyle{} << " or "; kU32Matcher.print(nullptr, out);}
+ kI32Matcher.print(nullptr, out); out << style::Plain(" or "); kU32Matcher.print(nullptr, out);}
 };
 
 /// EnumMatcher for 'match f32_texel_format'
 constexpr NumberMatcher kF32TexelFormatMatcher {
 /* match */ [](MatchState&, Number number) -> Number {
     switch (static_cast<core::TexelFormat>(number.Value())) {
+      case core::TexelFormat::kR8Unorm:
       case core::TexelFormat::kBgra8Unorm:
       case core::TexelFormat::kRgba8Unorm:
       case core::TexelFormat::kRgba8Snorm:
@@ -1542,7 +1543,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "bgra8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8unorm"<< TextStyle{} << ", " << style::Enum << "rgba8snorm"<< TextStyle{} << ", " << style::Enum << "rgba16float"<< TextStyle{} << ", " << style::Enum << "r32float"<< TextStyle{} << ", " << style::Enum << "rg32float"<< TextStyle{} << " or " << style::Enum << "rgba32float";
+  out<< style::Enum("r8unorm")<< style::Plain(", ") << style::Enum("bgra8unorm")<< style::Plain(", ") << style::Enum("rgba8unorm")<< style::Plain(", ") << style::Enum("rgba8snorm")<< style::Plain(", ") << style::Enum("rgba16float")<< style::Plain(", ") << style::Enum("r32float")<< style::Plain(", ") << style::Enum("rg32float")<< style::Plain(" or ") << style::Enum("rgba32float");
   }
 };
 
@@ -1561,7 +1562,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8sint"<< TextStyle{} << ", " << style::Enum << "rgba16sint"<< TextStyle{} << ", " << style::Enum << "r32sint"<< TextStyle{} << ", " << style::Enum << "rg32sint"<< TextStyle{} << " or " << style::Enum << "rgba32sint";
+  out<< style::Enum("rgba8sint")<< style::Plain(", ") << style::Enum("rgba16sint")<< style::Plain(", ") << style::Enum("r32sint")<< style::Plain(", ") << style::Enum("rg32sint")<< style::Plain(" or ") << style::Enum("rgba32sint");
   }
 };
 
@@ -1580,7 +1581,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "rgba8uint"<< TextStyle{} << ", " << style::Enum << "rgba16uint"<< TextStyle{} << ", " << style::Enum << "r32uint"<< TextStyle{} << ", " << style::Enum << "rg32uint"<< TextStyle{} << " or " << style::Enum << "rgba32uint";
+  out<< style::Enum("rgba8uint")<< style::Plain(", ") << style::Enum("rgba16uint")<< style::Plain(", ") << style::Enum("r32uint")<< style::Plain(", ") << style::Enum("rg32uint")<< style::Plain(" or ") << style::Enum("rgba32uint");
   }
 };
 
@@ -1593,7 +1594,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "write";
+  out<< style::Enum("write");
   }
 };
 
@@ -1606,7 +1607,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read_write";
+  out<< style::Enum("read_write");
   }
 };
 
@@ -1622,7 +1623,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "read"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("read")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
@@ -1638,7 +1639,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "write"<< TextStyle{} << " or " << style::Enum << "read_write";
+  out<< style::Enum("write")<< style::Plain(" or ") << style::Enum("read_write");
   }
 };
 
@@ -1655,7 +1656,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "function"<< TextStyle{} << ", " << style::Enum << "private"<< TextStyle{} << " or " << style::Enum << "workgroup";
+  out<< style::Enum("function")<< style::Plain(", ") << style::Enum("private")<< style::Plain(" or ") << style::Enum("workgroup");
   }
 };
 
@@ -1671,7 +1672,7 @@
     }
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "workgroup"<< TextStyle{} << " or " << style::Enum << "storage";
+  out<< style::Enum("workgroup")<< style::Plain(" or ") << style::Enum("storage");
   }
 };
 
@@ -1684,7 +1685,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "storage";
+  out<< style::Enum("storage");
   }
 };
 
@@ -1697,7 +1698,7 @@
     return Number::invalid;
   },
 /* print */ [](MatchState*, StyledText& out) {
-  out<< style::Enum << "workgroup";
+  out<< style::Enum("workgroup");
   }
 };
 
diff --git a/src/tint/lang/wgsl/resolver/BUILD.bazel b/src/tint/lang/wgsl/resolver/BUILD.bazel
index 51904b9..9893abb 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.bazel
+++ b/src/tint/lang/wgsl/resolver/BUILD.bazel
@@ -121,6 +121,7 @@
     "f16_extension_test.cc",
     "framebuffer_fetch_extension_test.cc",
     "function_validation_test.cc",
+    "graphite_extension_test.cc",
     "host_shareable_validation_test.cc",
     "increment_decrement_validation_test.cc",
     "inferred_type_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/BUILD.cmake b/src/tint/lang/wgsl/resolver/BUILD.cmake
index 6cb53ed..429b9c0 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.cmake
+++ b/src/tint/lang/wgsl/resolver/BUILD.cmake
@@ -119,6 +119,7 @@
   lang/wgsl/resolver/f16_extension_test.cc
   lang/wgsl/resolver/framebuffer_fetch_extension_test.cc
   lang/wgsl/resolver/function_validation_test.cc
+  lang/wgsl/resolver/graphite_extension_test.cc
   lang/wgsl/resolver/host_shareable_validation_test.cc
   lang/wgsl/resolver/increment_decrement_validation_test.cc
   lang/wgsl/resolver/inferred_type_test.cc
diff --git a/src/tint/lang/wgsl/resolver/BUILD.gn b/src/tint/lang/wgsl/resolver/BUILD.gn
index 18cab8b..9456dbb 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.gn
+++ b/src/tint/lang/wgsl/resolver/BUILD.gn
@@ -121,6 +121,7 @@
       "f16_extension_test.cc",
       "framebuffer_fetch_extension_test.cc",
       "function_validation_test.cc",
+      "graphite_extension_test.cc",
       "host_shareable_validation_test.cc",
       "increment_decrement_validation_test.cc",
       "inferred_type_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/graphite_extension_test.cc b/src/tint/lang/wgsl/resolver/graphite_extension_test.cc
new file mode 100644
index 0000000..71f727e
--- /dev/null
+++ b/src/tint/lang/wgsl/resolver/graphite_extension_test.cc
@@ -0,0 +1,72 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/access.h"
+#include "src/tint/lang/core/texel_format.h"
+#include "src/tint/lang/core/type/texture_dimension.h"
+#include "src/tint/lang/wgsl/resolver/resolver.h"
+#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
+
+#include "gmock/gmock.h"
+
+namespace tint::resolver {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using ResolverGraphiteExtensionTest = ResolverTest;
+
+TEST_F(ResolverGraphiteExtensionTest, r8unorm_UseWithoutExtension) {
+    // @group(0) @binding(0) var T: texture_storage_2d<r8unorm, write>;
+
+    GlobalVar("T",
+              ty.storage_texture(core::type::TextureDimension::k2d, core::TexelFormat::kR8Unorm,
+                                 core::Access::kWrite),
+              Vector{Group(0_a), Binding(0_a)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(error: 'r8unorm' requires the 'chromium_internal_graphite' extension)");
+}
+
+TEST_F(ResolverGraphiteExtensionTest, r8unorm_UseWithExtension) {
+    // enable chromium_internal_graphite;
+    // @group(0) @binding(0) var T: texture_storage_3d<r8unorm, read_write>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalGraphite);
+
+    GlobalVar("T",
+              ty.storage_texture(core::type::TextureDimension::k2d, core::TexelFormat::kR8Unorm,
+                                 core::Access::kWrite),
+              Vector{Group(0_a), Binding(0_a)});
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+}  // namespace
+}  // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 20ec2a4..b1a1775 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -31,11 +31,13 @@
 #include <cmath>
 #include <iomanip>
 #include <limits>
+#include <string_view>
 #include <utility>
 
 #include "src/tint/lang/core/builtin_type.h"
 #include "src/tint/lang/core/constant/scalar.h"
 #include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/core/texel_format.h"
 #include "src/tint/lang/core/type/abstract_float.h"
 #include "src/tint/lang/core/type/abstract_int.h"
 #include "src/tint/lang/core/type/array.h"
@@ -271,8 +273,8 @@
             attribute,  //
             [&](const ast::InternalAttribute* attr) -> bool { return InternalAttribute(attr); },
             [&](Default) {
-                ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "let"
-                                                              << style::Plain << " declaration");
+                ErrorInvalidAttribute(attribute,
+                                      StyledText{} << style::Keyword("let") << " declaration");
                 return false;
             });
         if (!ok) {
@@ -281,8 +283,7 @@
     }
 
     if (TINT_UNLIKELY(!v->initializer)) {
-        AddError(v->source) << style::Keyword << "let" << style::Plain
-                            << " declaration must have an initializer";
+        AddError(v->source) << style::Keyword("let") << " declaration must have an initializer";
         return nullptr;
     }
 
@@ -303,8 +304,8 @@
 
     if (!ApplyAddressSpaceUsageToType(core::AddressSpace::kUndefined,
                                       const_cast<core::type::Type*>(sem->Type()), v->source)) {
-        AddNote(v->source) << "while instantiating " << style::Keyword << "let " << style::Variable
-                           << v->name->symbol.Name();
+        AddNote(v->source) << "while instantiating " << style::Keyword("let ")
+                           << style::Variable(v->name->symbol.NameView());
         return nullptr;
     }
 
@@ -362,8 +363,8 @@
 
     if (!ApplyAddressSpaceUsageToType(core::AddressSpace::kUndefined,
                                       const_cast<core::type::Type*>(ty), v->source)) {
-        AddNote(v->source) << "while instantiating " << style::Keyword << "override "
-                           << style::Variable << v->name->symbol.Name();
+        AddNote(v->source) << "while instantiating " << style::Keyword("override ")
+                           << style::Variable(v->name->symbol.NameView());
         return nullptr;
     }
 
@@ -381,23 +382,21 @@
                 }
                 if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
                     AddError(attr->source)
-                        << style::Attribute << "@id" << style::Plain << " must be an "
-                        << style::Type << "i32" << style::Plain << " or " << style::Type << "u32"
-                        << style::Plain << " value";
+                        << style::Attribute("@id") << " must be an " << style::Type("i32") << " or "
+                        << style::Type("u32") << " value";
                     return false;
                 }
 
                 auto const_value = materialized->ConstantValue();
                 auto value = const_value->ValueAs<AInt>();
                 if (value < 0) {
-                    AddError(attr->source) << style::Attribute << "@id" << style::Plain
-                                           << " value must be non-negative";
+                    AddError(attr->source)
+                        << style::Attribute("@id") << " value must be non-negative";
                     return false;
                 }
                 if (value > std::numeric_limits<decltype(OverrideId::value)>::max()) {
                     AddError(attr->source)
-                        << style::Attribute << "@id" << style::Plain
-                        << " value must be between 0 and "
+                        << style::Attribute("@id") << " value must be between 0 and "
                         << std::numeric_limits<decltype(OverrideId::value)>::max();
                     return false;
                 }
@@ -410,8 +409,8 @@
                 return true;
             },
             [&](Default) {
-                ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "override"
-                                                              << style::Plain << " declaration");
+                ErrorInvalidAttribute(attribute,
+                                      StyledText{} << style::Keyword("override") << " declaration");
                 return false;
             });
         if (!ok) {
@@ -435,13 +434,13 @@
 
     for (auto* attribute : c->attributes) {
         Mark(attribute);
-        bool ok = Switch(attribute,  //
-                         [&](Default) {
-                             ErrorInvalidAttribute(attribute,
-                                                   StyledText{} << style::Keyword << "const"
-                                                                << style::Plain << " declaration");
-                             return false;
-                         });
+        bool ok =
+            Switch(attribute,  //
+                   [&](Default) {
+                       ErrorInvalidAttribute(
+                           attribute, StyledText{} << style::Keyword("const") << " declaration");
+                       return false;
+                   });
         if (!ok) {
             return nullptr;
         }
@@ -493,7 +492,7 @@
 
     if (!ApplyAddressSpaceUsageToType(core::AddressSpace::kUndefined,
                                       const_cast<core::type::Type*>(ty), c->source)) {
-        AddNote(c->source) << "while instantiating 'const' " << c->name->symbol.Name();
+        AddNote(c->source) << "while instantiating 'const' " << c->name->symbol.NameView();
         return nullptr;
     }
 
@@ -606,7 +605,7 @@
     if (!ApplyAddressSpaceUsageToType(sem->AddressSpace(),
                                       const_cast<core::type::Type*>(sem->Type()),
                                       var->type ? var->type->source : var->source)) {
-        AddNote(var->source) << "while instantiating 'var' " << var->name->symbol.Name();
+        AddNote(var->source) << "while instantiating 'var' " << var->name->symbol.NameView();
         return nullptr;
     }
 
@@ -698,8 +697,8 @@
                 case kErrored:
                     return nullptr;
                 case kInvalid:
-                    ErrorInvalidAttribute(
-                        attribute, StyledText{} << "module-scope " << style::Keyword << "var");
+                    ErrorInvalidAttribute(attribute,
+                                          StyledText{} << "module-scope " << style::Keyword("var"));
                     return nullptr;
             }
         }
@@ -716,7 +715,7 @@
                 [&](const ast::InternalAttribute* attr) { return InternalAttribute(attr); },
                 [&](Default) {
                     ErrorInvalidAttribute(
-                        attribute, StyledText{} << "function-scope " << style::Keyword << "var");
+                        attribute, StyledText{} << "function-scope " << style::Keyword("var"));
                     return false;
                 });
             if (!ok) {
@@ -737,7 +736,8 @@
     b.Sem().Add(param, sem);
 
     auto add_note = [&] {
-        AddNote(param->source) << "while instantiating parameter " << param->name->symbol.Name();
+        AddNote(param->source) << "while instantiating parameter "
+                               << param->name->symbol.NameView();
     };
 
     if (func->IsEntryPoint()) {
@@ -1032,7 +1032,7 @@
 
         {  // Check the parameter name is unique for the function
             if (auto added = parameter_names.Add(param->name->symbol, param->source); !added) {
-                auto name = param->name->symbol.Name();
+                auto name = param->name->symbol.NameView();
                 AddError(param->source) << "redefinition of parameter '" << name << "'";
                 AddNote(added.value) << "previous definition is here";
                 return nullptr;
@@ -1158,7 +1158,7 @@
     if (auto* str = return_type->As<core::type::Struct>()) {
         if (!ApplyAddressSpaceUsageToType(core::AddressSpace::kUndefined, str, decl->source)) {
             AddNote(decl->source) << "while instantiating return type for "
-                                  << decl->name->symbol.Name();
+                                  << decl->name->symbol.NameView();
             return nullptr;
         }
 
@@ -1725,7 +1725,7 @@
                 break;
             case Alias::ModuleScope: {
                 auto* func = var.expr->Stmt()->Function();
-                auto func_name = func->Declaration()->name->symbol.Name();
+                auto func_name = func->Declaration()->name->symbol.NameView();
                 AddNote(var.expr->Declaration()->source)
                     << "aliases with module-scope variable " << var.access << " in '" << func_name
                     << "'";
@@ -2721,7 +2721,7 @@
             break;
     }
 
-    auto name = ident->symbol.Name();
+    auto name = ident->symbol.NameView();
     StringStream err;
     err << " unhandled builtin type '" << name << "'";
     AddICE(err.str(), ident->source);
@@ -3004,9 +3004,8 @@
     auto* tmpl_ident = ident->As<ast::TemplatedIdentifier>();
     if (!tmpl_ident) {
         if (TINT_UNLIKELY(min_args != 0)) {
-            AddError(Source{ident->source.range.end})
-                << "expected " << style::Code << "<" << style::Plain << " for " << style::Code
-                << ident->symbol.Name();
+            AddError(Source{ident->source.range.end}) << "expected " << style::Code("<") << " for "
+                                                      << style::Code(ident->symbol.NameView());
         }
         return nullptr;
     }
@@ -3021,19 +3020,19 @@
     }
     if (min_args == max_args) {
         if (TINT_UNLIKELY(ident->arguments.Length() != min_args)) {
-            AddError(ident->source) << style::Code << ident->symbol.Name() << style::Plain
-                                    << " requires " << min_args << " template arguments";
+            AddError(ident->source) << style::Code(ident->symbol.NameView()) << " requires "
+                                    << min_args << " template arguments";
             return false;
         }
     } else {
         if (TINT_UNLIKELY(ident->arguments.Length() < min_args)) {
-            AddError(ident->source) << style::Code << ident->symbol.Name() << style::Plain
+            AddError(ident->source) << style::Code(ident->symbol.NameView())
                                     << " requires at least " << min_args << " template arguments";
             return false;
         }
         if (TINT_UNLIKELY(ident->arguments.Length() > max_args)) {
-            AddError(ident->source) << style::Code << ident->symbol.Name() << style::Plain
-                                    << " requires at most " << max_args << " template arguments";
+            AddError(ident->source) << style::Code(ident->symbol.NameView()) << " requires at most "
+                                    << max_args << " template arguments";
             return false;
         }
     }
@@ -3257,7 +3256,7 @@
     auto resolved = dependencies_.resolved_identifiers.Get(ident);
     if (!resolved) {
         StringStream err;
-        err << "identifier '" << ident->symbol.Name() << "' was not resolved";
+        err << "identifier '" << ident->symbol.NameView() << "' was not resolved";
         AddICE(err.str(), expr->source);
         return nullptr;
     }
@@ -3299,11 +3298,11 @@
                                 if (decl->order >= loop_block->NumDeclsAtFirstContinue()) {
                                     AddError(loop_block->FirstContinue()->source)
                                         << "continue statement bypasses declaration of '"
-                                        << symbol.Name() << "'";
+                                        << symbol.NameView() << "'";
                                     AddNote(decl->variable->Declaration()->source)
-                                        << "identifier '" << symbol.Name() << "' declared here";
+                                        << "identifier '" << symbol.NameView() << "' declared here";
                                     AddNote(expr->source)
-                                        << "identifier '" << symbol.Name()
+                                        << "identifier '" << symbol.NameView()
                                         << "' referenced in continuing block here";
                                     return nullptr;
                                 }
@@ -3319,8 +3318,8 @@
                     if (!current_function_ && variable->Declaration()->Is<ast::Var>()) {
                         // Use of a module-scope 'var' outside of a function.
                         AddError(expr->source)
-                            << style::Keyword << "var " << style::Variable << ident->symbol.Name()
-                            << style::Plain << " cannot be referenced at module-scope";
+                            << style::Keyword("var ") << style::Variable(ident->symbol.NameView())
+                            << " cannot be referenced at module-scope";
                         sem_.NoteDeclarationSource(variable->Declaration());
                         return nullptr;
                     }
@@ -3462,7 +3461,7 @@
             }
 
             if (member == nullptr) {
-                AddError(expr->source) << "struct member " << symbol.Name() << " not found";
+                AddError(expr->source) << "struct member " << symbol.NameView() << " not found";
                 return nullptr;
             }
 
@@ -3483,7 +3482,7 @@
         },
 
         [&](const core::type::Vector* vec) -> sem::ValueExpression* {
-            std::string s = expr->member->symbol.Name();
+            std::string_view s = expr->member->symbol.NameView();
             auto size = s.size();
             Vector<uint32_t, 4> swizzle;
             swizzle.Reserve(s.size());
@@ -3707,7 +3706,7 @@
                 root_ident = expr->RootIdentifier();
             } else {
                 AddError(unary->expr->source) << "cannot dereference expression of type "
-                                              << style::Type << sem_.TypeNameOf(expr_ty);
+                                              << style::Type(sem_.TypeNameOf(expr_ty));
                 return nullptr;
             }
             break;
@@ -3767,17 +3766,15 @@
     }
 
     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-        AddError(attr->source) << style::Attribute << "@location" << style::Plain << " must be an "
-                               << style::Type << "i32" << style::Plain << " or " << style::Type
-                               << "u32" << style::Plain << " value";
+        AddError(attr->source) << style::Attribute("@location") << " must be an "
+                               << style::Type("i32") << " or " << style::Type("u32") << " value";
         return Failure{};
     }
 
     auto const_value = materialized->ConstantValue();
     auto value = const_value->ValueAs<AInt>();
     if (value < 0) {
-        AddError(attr->source) << style::Attribute << "@location" << style::Plain
-                               << " value must be non-negative";
+        AddError(attr->source) << style::Attribute("@location") << " value must be non-negative";
         return Failure{};
     }
 
@@ -3794,17 +3791,15 @@
     }
 
     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-        AddError(attr->source) << style::Attribute << "@color" << style::Plain << " must be an "
-                               << style::Type << "i32" << style::Plain << " or " << style::Type
-                               << "u32" << style::Plain << " value";
+        AddError(attr->source) << style::Attribute("@color") << " must be an " << style::Type("i32")
+                               << " or " << style::Type("u32") << " value";
         return Failure{};
     }
 
     auto const_value = materialized->ConstantValue();
     auto value = const_value->ValueAs<AInt>();
     if (value < 0) {
-        AddError(attr->source) << style::Attribute << "@color" << style::Plain
-                               << " value must be non-negative";
+        AddError(attr->source) << style::Attribute("@color") << " value must be non-negative";
         return Failure{};
     }
 
@@ -3820,17 +3815,15 @@
     }
 
     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-        AddError(attr->source) << style::Attribute << "@blend_src" << style::Plain << style::Type
-                               << "i32" << style::Plain << " or " << style::Type << "u32"
-                               << style::Plain << " value";
+        AddError(attr->source) << style::Attribute("@blend_src") << style::Type("i32") << " or "
+                               << style::Type("u32") << " value";
         return Failure{};
     }
 
     auto const_value = materialized->ConstantValue();
     auto value = const_value->ValueAs<AInt>();
     if (value != 0 && value != 1) {
-        AddError(attr->source) << style::Attribute << "@blend_src" << style::Plain
-                               << " value must be zero or one";
+        AddError(attr->source) << style::Attribute("@blend_src") << " value must be zero or one";
         return Failure{};
     }
 
@@ -3846,17 +3839,15 @@
         return Failure{};
     }
     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-        AddError(attr->source) << style::Attribute << "@binding" << style::Plain << " must be an "
-                               << style::Type << "i32" << style::Plain << " or " << style::Type
-                               << "u32" << style::Plain << " value";
+        AddError(attr->source) << style::Attribute("@binding") << " must be an "
+                               << style::Type("i32") << " or " << style::Type("u32") << " value";
         return Failure{};
     }
 
     auto const_value = materialized->ConstantValue();
     auto value = const_value->ValueAs<AInt>();
     if (value < 0) {
-        AddError(attr->source) << style::Attribute << "@binding" << style::Plain
-                               << " value must be non-negative";
+        AddError(attr->source) << style::Attribute("@binding") << " value must be non-negative";
         return Failure{};
     }
     return static_cast<uint32_t>(value);
@@ -3871,17 +3862,15 @@
         return Failure{};
     }
     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-        AddError(attr->source) << style::Attribute << "@group" << style::Plain << " must be an "
-                               << style::Type << "i32" << style::Plain << " or " << style::Type
-                               << "u32" << style::Plain << " value";
+        AddError(attr->source) << style::Attribute("@group") << " must be an " << style::Type("i32")
+                               << " or " << style::Type("u32") << " value";
         return Failure{};
     }
 
     auto const_value = materialized->ConstantValue();
     auto value = const_value->ValueAs<AInt>();
     if (value < 0) {
-        AddError(attr->source) << style::Attribute << "@group" << style::Plain
-                               << " value must be non-negative";
+        AddError(attr->source) << style::Attribute("@group") << " value must be non-negative";
         return Failure{};
     }
     return static_cast<uint32_t>(value);
@@ -3899,11 +3888,10 @@
     Vector<const core::type::Type*, 3> arg_tys;
 
     auto err_bad_expr = [&](const ast::Expression* value) {
-        AddError(value->source) << style::Attribute << "@workgroup_size" << style::Plain
+        AddError(value->source) << style::Attribute("@workgroup_size")
                                 << " argument must be a constant or override-expression of type "
-                                << style::Type << "abstract-integer" << style::Plain << ", "
-                                << style::Type << "i32" << style::Plain << " or " << style::Type
-                                << "u32";
+                                << style::Type("abstract-integer") << ", " << style::Type("i32")
+                                << " or " << style::Type("u32");
     };
 
     for (size_t i = 0; i < 3; i++) {
@@ -3935,9 +3923,9 @@
 
     auto* common_ty = core::type::Type::Common(arg_tys);
     if (!common_ty) {
-        AddError(attr->source) << style::Attribute << "@workgroup_size" << style::Plain
-                               << " arguments must be of the same type, either " << style::Type
-                               << "i32" << style::Plain << " or " << style::Type << "u32";
+        AddError(attr->source) << style::Attribute("@workgroup_size")
+                               << " arguments must be of the same type, either "
+                               << style::Type("i32") << " or " << style::Type("u32");
         return Failure{};
     }
 
@@ -3953,8 +3941,8 @@
         }
         if (auto* value = materialized->ConstantValue()) {
             if (value->ValueAs<AInt>() < 1) {
-                AddError(values[i]->source) << style::Attribute << "@workgroup_size" << style::Plain
-                                            << " argument must be at least 1";
+                AddError(values[i]->source)
+                    << style::Attribute("@workgroup_size") << " argument must be at least 1";
                 return Failure{};
             }
             ws[i] = value->ValueAs<u32>();
@@ -4037,18 +4025,18 @@
 bool Resolver::DiagnosticControl(const ast::DiagnosticControl& control) {
     Mark(control.rule_name);
     Mark(control.rule_name->name);
-    auto name = control.rule_name->name->symbol.Name();
+    auto name = control.rule_name->name->symbol.NameView();
 
     if (control.rule_name->category) {
         Mark(control.rule_name->category);
-        if (control.rule_name->category->symbol.Name() == "chromium") {
+        if (control.rule_name->category->symbol.NameView() == "chromium") {
             auto rule = wgsl::ParseChromiumDiagnosticRule(name);
             if (rule != wgsl::ChromiumDiagnosticRule::kUndefined) {
                 validator_.DiagnosticFilters().Set(rule, control.severity);
             } else {
                 auto& warning = AddWarning(control.rule_name->source)
-                                << "unrecognized diagnostic rule " << style::Code << "chromium."
-                                << name << style::Plain << "\n";
+                                << "unrecognized diagnostic rule " << style::Code("chromium.", name)
+                                << "\n";
                 tint::SuggestAlternativeOptions opts;
                 opts.prefix = "chromium.";
                 tint::SuggestAlternatives(name, wgsl::kChromiumDiagnosticRuleStrings,
@@ -4063,8 +4051,7 @@
         validator_.DiagnosticFilters().Set(rule, control.severity);
     } else {
         auto& warning = AddWarning(control.rule_name->source)
-                        << "unrecognized diagnostic rule " << style::Code << name << style::Plain
-                        << "\n";
+                        << "unrecognized diagnostic rule " << style::Code(name) << "\n";
         tint::SuggestAlternatives(name, wgsl::kCoreDiagnosticRuleStrings, warning.message);
     }
     return true;
@@ -4075,7 +4062,7 @@
         Mark(ext);
         enabled_extensions_.Add(ext->name);
         if (!allowed_features_.extensions.count(ext->name)) {
-            AddError(ext->source) << "extension " << style::Code << ext->name << style::Plain
+            AddError(ext->source) << "extension " << style::Code(ext->name)
                                   << " is not allowed in the current environment";
             return false;
         }
@@ -4086,8 +4073,8 @@
 bool Resolver::Requires(const ast::Requires* req) {
     for (auto feature : req->features) {
         if (!allowed_features_.features.count(feature)) {
-            AddError(req->source) << "language feature " << style::Code << wgsl::ToString(feature)
-                                  << style::Plain << " is not allowed in the current environment";
+            AddError(req->source) << "language feature " << style::Code(wgsl::ToString(feature))
+                                  << " is not allowed in the current environment";
             return false;
         }
     }
@@ -4144,7 +4131,7 @@
             if (auto* ty = count_val->Type(); !ty->is_integer_scalar()) {
                 AddError(count_expr->source)
                     << "array count must evaluate to a constant integer expression, but is type "
-                    << style::Type << ty->FriendlyName();
+                    << style::Type(ty->FriendlyName());
                 return nullptr;
             }
 
@@ -4192,8 +4179,7 @@
                 return true;
             },
             [&](Default) {
-                ErrorInvalidAttribute(
-                    attribute, StyledText{} << style::Type << "array" << style::Plain << " types");
+                ErrorInvalidAttribute(attribute, StyledText{} << style::Type("array") << " types");
                 return false;
             });
         if (!ok) {
@@ -4260,7 +4246,7 @@
 
 sem::Struct* Resolver::Structure(const ast::Struct* str) {
     auto struct_name = [&] {  //
-        return str->name->symbol.Name();
+        return str->name->symbol.NameView();
     };
 
     if (validator_.IsValidationEnabled(str->attributes,
@@ -4269,9 +4255,9 @@
         // https://gpuweb.github.io/gpuweb/wgsl/#limits
         const size_t kMaxNumStructMembers = 16383;
         if (str->members.Length() > kMaxNumStructMembers) {
-            AddError(str->source) << style::Keyword << "struct " << style::Type << struct_name()
-                                  << style::Plain << " has " << str->members.Length()
-                                  << " members, maximum is " << kMaxNumStructMembers;
+            AddError(str->source) << style::Keyword("struct ") << style::Type(struct_name())
+                                  << " has " << str->members.Length() << " members, maximum is "
+                                  << kMaxNumStructMembers;
             return nullptr;
         }
     }
@@ -4285,8 +4271,8 @@
         bool ok = Switch(
             attribute, [&](const ast::InternalAttribute* attr) { return InternalAttribute(attr); },
             [&](Default) {
-                ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "struct"
-                                                              << style::Plain << " declarations");
+                ErrorInvalidAttribute(attribute,
+                                      StyledText{} << style::Keyword("struct") << " declarations");
                 return false;
             });
         if (!ok) {
@@ -4313,7 +4299,7 @@
         Mark(member->name);
         if (auto added = member_map.Add(member->name->symbol, member); !added) {
             AddError(member->source)
-                << "redefinition of " << style::Code << member->name->symbol.Name();
+                << "redefinition of " << style::Code(member->name->symbol.NameView());
             AddNote(added.value->source) << "previous definition is here";
             return nullptr;
         }
@@ -4384,22 +4370,22 @@
                         return false;
                     }
                     if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
-                        AddError(attr->source) << style::Attribute << "@align" << style::Plain
-                                               << " value must be an " << style::Type << "i32"
-                                               << style::Plain << " or " << style::Type << "u32";
+                        AddError(attr->source)
+                            << style::Attribute("@align") << " value must be an "
+                            << style::Type("i32") << " or " << style::Type("u32");
                         return false;
                     }
 
                     auto const_value = materialized->ConstantValue();
                     if (!const_value) {
-                        AddError(attr->source) << style::Attribute << "@align" << style::Plain
-                                               << " value must be constant expression";
+                        AddError(attr->source)
+                            << style::Attribute("@align") << " value must be constant expression";
                         return false;
                     }
                     auto value = const_value->ValueAs<AInt>();
 
                     if (value <= 0 || !tint::IsPowerOfTwo(value)) {
-                        AddError(attr->source) << style::Attribute << "@align" << style::Plain
+                        AddError(attr->source) << style::Attribute("@align")
                                                << " value must be a positive, power-of-two integer";
                         return false;
                     }
@@ -4416,30 +4402,30 @@
                         return false;
                     }
                     if (!materialized->Type()->IsAnyOf<core::type::U32, core::type::I32>()) {
-                        AddError(attr->source) << style::Attribute << "@size" << style::Plain
-                                               << " value must be an " << style::Type << "i32"
-                                               << style::Plain << " or " << style::Type << "u32";
+                        AddError(attr->source)
+                            << style::Attribute("@size") << " value must be an "
+                            << style::Type("i32") << " or " << style::Type("u32");
                         return false;
                     }
 
                     auto const_value = materialized->ConstantValue();
                     if (!const_value) {
-                        AddError(attr->expr->source) << style::Attribute << "@size" << style::Plain
-                                                     << " value must be constant expression";
+                        AddError(attr->expr->source)
+                            << style::Attribute("@size") << " value must be constant expression";
                         return false;
                     }
                     {
                         auto value = const_value->ValueAs<AInt>();
                         if (value <= 0) {
-                            AddError(attr->source) << style::Attribute << "@size" << style::Plain
-                                                   << " value must be a positive integer";
+                            AddError(attr->source)
+                                << style::Attribute("@size") << " value must be a positive integer";
                             return false;
                         }
                     }
                     auto value = const_value->ValueAs<uint64_t>();
                     if (value < size) {
                         AddError(attr->source)
-                            << style::Attribute << "@size" << style::Plain
+                            << style::Attribute("@size")
                             << " must be at least as big as the type's size (" << size << ")";
                         return false;
                     }
@@ -4497,17 +4483,16 @@
                 [&](const ast::StrideAttribute* attr) {
                     if (validator_.IsValidationEnabled(
                             member->attributes, ast::DisabledValidation::kIgnoreStrideAttribute)) {
-                        ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "struct"
-                                                                      << style::Plain
-                                                                      << " members");
+                        ErrorInvalidAttribute(
+                            attribute, StyledText{} << style::Keyword("struct") << " members");
                         return false;
                     }
                     return StrideAttribute(attr);
                 },
                 [&](const ast::InternalAttribute* attr) { return InternalAttribute(attr); },
                 [&](Default) {
-                    ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "struct"
-                                                                  << style::Plain << " members");
+                    ErrorInvalidAttribute(attribute,
+                                          StyledText{} << style::Keyword("struct") << " members");
                     return false;
                 });
             if (!ok) {
@@ -4516,9 +4501,9 @@
         }
 
         if (has_offset_attr && (has_align_attr || has_size_attr)) {
-            AddError(member->source) << style::Attribute << "@offset" << style::Plain
-                                     << " cannot be used with " << style::Attribute << "@align"
-                                     << style::Plain << " or " << style::Attribute << "@size";
+            AddError(member->source)
+                << style::Attribute("@offset") << " cannot be used with "
+                << style::Attribute("@align") << " or " << style::Attribute("@size");
             return nullptr;
         }
 
@@ -4583,9 +4568,9 @@
     //  https://gpuweb.github.io/gpuweb/wgsl/#limits
     const size_t nest_depth = 1 + members_nest_depth;
     if (nest_depth > kMaxNestDepthOfCompositeType) {
-        AddError(str->source) << style::Keyword << "struct " << style::Type << struct_name()
-                              << style::Plain << " has nesting depth of " << nest_depth
-                              << ", maximum is " << kMaxNestDepthOfCompositeType;
+        AddError(str->source) << style::Keyword("struct ") << style::Type(struct_name())
+                              << " has nesting depth of " << nest_depth << ", maximum is "
+                              << kMaxNestDepthOfCompositeType;
         return nullptr;
     }
     nest_depth_.Add(out, nest_depth);
@@ -4673,8 +4658,8 @@
                 attribute,
                 [&](const ast::DiagnosticAttribute* attr) { return DiagnosticAttribute(attr); },
                 [&](Default) {
-                    ErrorInvalidAttribute(attribute, StyledText{} << style::Keyword << "switch"
-                                                                  << style::Plain << " body");
+                    ErrorInvalidAttribute(attribute,
+                                          StyledText{} << style::Keyword("switch") << " body");
                     return false;
                 });
             if (!ok) {
@@ -4930,9 +4915,9 @@
     }
 
     if (core::IsHostShareable(address_space) && !validator_.IsHostShareable(ty)) {
-        AddError(usage) << "type " << style::Type << sem_.TypeNameOf(ty) << style::Plain
-                        << " cannot be used in address space " << style::Enum << address_space
-                        << style::Plain << " as it is non-host-shareable";
+        AddError(usage) << "type " << style::Type(sem_.TypeNameOf(ty))
+                        << " cannot be used in address space " << style::Enum(address_space)
+                        << " as it is non-host-shareable";
         return false;
     }
 
@@ -5040,7 +5025,7 @@
 
 bool Resolver::CheckNotTemplated(const char* use, const ast::Identifier* ident) {
     if (TINT_UNLIKELY(ident->Is<ast::TemplatedIdentifier>())) {
-        AddError(ident->source) << use << " " << style::Code << ident->symbol.Name() << style::Plain
+        AddError(ident->source) << use << " " << style::Code(ident->symbol.NameView())
                                 << " does not take template arguments";
         if (auto resolved = dependencies_.resolved_identifiers.Get(ident)) {
             if (auto* ast_node = resolved->Node()) {
@@ -5053,8 +5038,7 @@
 }
 
 void Resolver::ErrorInvalidAttribute(const ast::Attribute* attr, StyledText use) {
-    AddError(attr->source) << style::Attribute << "@" << attr->Name() << style::Plain
-                           << " is not valid for " << use;
+    AddError(attr->source) << style::Attribute("@", attr->Name()) << " is not valid for " << use;
 }
 
 void Resolver::AddICE(std::string_view msg, const Source& source) const {
diff --git a/src/tint/lang/wgsl/resolver/sem_helper.cc b/src/tint/lang/wgsl/resolver/sem_helper.cc
index 5164284..494e4f8 100644
--- a/src/tint/lang/wgsl/resolver/sem_helper.cc
+++ b/src/tint/lang/wgsl/resolver/sem_helper.cc
@@ -71,8 +71,7 @@
     auto* type = ty_expr->Type();
     if (auto* incomplete = type->As<IncompleteType>(); TINT_UNLIKELY(incomplete)) {
         AddError(expr->Declaration()->source.End())
-            << "expected " << style::Code << "<" << style::Plain << " for " << style::Type
-            << incomplete->builtin;
+            << "expected " << style::Code("<") << " for " << style::Type(incomplete->builtin);
         return nullptr;
     }
 
@@ -86,55 +85,57 @@
         expr,  //
         [&](const sem::VariableUser* var_expr) {
             auto* variable = var_expr->Variable()->Declaration();
-            auto name = variable->name->symbol.Name();
+            auto name = variable->name->symbol.NameView();
             Switch(
-                variable,                                                             //
-                [&](const ast::Var*) { text << style::Keyword << "var"; },            //
-                [&](const ast::Let*) { text << style::Keyword << "let"; },            //
-                [&](const ast::Const*) { text << style::Keyword << "const"; },        //
-                [&](const ast::Parameter*) { text << "parameter"; },                  //
-                [&](const ast::Override*) { text << style::Keyword << "override"; },  //
-                [&](Default) { text << "variable"; });
-            text << " " << style::Variable << name;
+                variable,                                                                         //
+                [&](const ast::Var*) { text << style::Keyword("var") << style::Code(" "); },      //
+                [&](const ast::Let*) { text << style::Keyword("let") << style::Code(" "); },      //
+                [&](const ast::Const*) { text << style::Keyword("const") << style::Code(" "); },  //
+                [&](const ast::Parameter*) { text << "parameter "; },                             //
+                [&](const ast::Override*) {
+                    text << style::Keyword("override") << style::Code(" ");
+                },  //
+                [&](Default) { text << "variable "; });
+            text << style::Variable(name);
         },
         [&](const sem::ValueExpression* val_expr) {
-            text << "value of type " << style::Type << val_expr->Type()->FriendlyName();
+            text << "value of type " << style::Type(val_expr->Type()->FriendlyName());
         },
         [&](const sem::TypeExpression* ty_expr) {
-            text << "type " << style::Type << ty_expr->Type()->FriendlyName();
+            text << "type " << style::Type(ty_expr->Type()->FriendlyName());
         },
         [&](const sem::FunctionExpression* fn_expr) {
             auto* fn = fn_expr->Function()->Declaration();
-            text << "function " << style::Function << fn->name->symbol.Name();
+            text << "function " << style::Function(fn->name->symbol.NameView());
         },
         [&](const sem::BuiltinEnumExpression<wgsl::BuiltinFn>* fn) {
-            text << "builtin function " << style::Function << fn->Value();
+            text << "builtin function " << style::Function(fn->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::Access>* access) {
-            text << "access " << style::Enum << access->Value();
+            text << "access " << style::Enum(access->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::AddressSpace>* addr) {
-            text << "address space " << style::Enum << addr->Value();
+            text << "address space " << style::Enum(addr->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::BuiltinValue>* builtin) {
-            text << "builtin value " << style::Enum << builtin->Value();
+            text << "builtin value " << style::Enum(builtin->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::InterpolationSampling>* fmt) {
-            text << "interpolation sampling " << style::Enum << fmt->Value();
+            text << "interpolation sampling " << style::Enum(fmt->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::InterpolationType>* fmt) {
-            text << "interpolation type " << style::Enum << fmt->Value();
+            text << "interpolation type " << style::Enum(fmt->Value());
         },
         [&](const sem::BuiltinEnumExpression<core::TexelFormat>* fmt) {
-            text << "texel format " << style::Enum << fmt->Value();
+            text << "texel format " << style::Enum(fmt->Value());
         },
         [&](const UnresolvedIdentifier* ui) {
-            auto name = ui->Identifier()->identifier->symbol.Name();
-            text << "unresolved identifier " << style::Code << name;
+            auto name = ui->Identifier()->identifier->symbol.NameView();
+            text << "unresolved identifier " << style::Code(name);
         },  //
         TINT_ICE_ON_NO_MATCH);
 
-    return text << style::Plain;
+    return text;
 }
 
 void SemHelper::ErrorUnexpectedExprKind(
@@ -143,8 +144,8 @@
     tint::Slice<const std::string_view> suggestions /* = Empty */) const {
     if (auto* ui = expr->As<UnresolvedIdentifier>()) {
         auto* ident = ui->Identifier();
-        auto name = ident->identifier->symbol.Name();
-        AddError(ident->source) << "unresolved " << wanted << " " << style::Code << name;
+        auto name = ident->identifier->symbol.NameView();
+        AddError(ident->source) << "unresolved " << wanted << " " << style::Code(name);
         if (!suggestions.IsEmpty()) {
             // Filter out suggestions that have a leading underscore.
             Vector<std::string_view, 8> filtered;
@@ -169,7 +170,7 @@
         if (expr->IsAnyOf<sem::FunctionExpression, sem::TypeExpression,
                           sem::BuiltinEnumExpression<wgsl::BuiltinFn>>()) {
             AddNote(ident->source.End())
-                << "are you missing " << style::Code << "()" << style::Plain << "?";
+                << "are you missing " << style::Code("()") << style::Plain("?");
         }
     }
 }
@@ -191,36 +192,36 @@
     Switch(
         node,
         [&](const ast::Struct* n) {
-            AddNote(n->source) << style::Keyword << "struct " << style::Type
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("struct ")
+                               << style::Type(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Alias* n) {
-            AddNote(n->source) << style::Keyword << "alias " << style::Type
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("alias ")
+                               << style::Type(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Var* n) {
-            AddNote(n->source) << style::Keyword << "var " << style::Variable
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("var ")
+                               << style::Variable(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Let* n) {
-            AddNote(n->source) << style::Keyword << "let " << style::Variable
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("let ")
+                               << style::Variable(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Override* n) {
-            AddNote(n->source) << style::Keyword << "override " << style::Variable
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("override ")
+                               << style::Variable(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Const* n) {
-            AddNote(n->source) << style::Keyword << "const " << style::Variable
-                               << n->name->symbol.Name() << style::Plain << " declared here";
+            AddNote(n->source) << style::Keyword("const ")
+                               << style::Variable(n->name->symbol.NameView()) << " declared here";
         },
         [&](const ast::Parameter* n) {
-            AddNote(n->source) << "parameter " << style::Variable << n->name->symbol.Name()
-                               << style::Plain << " declared here";
+            AddNote(n->source) << "parameter " << style::Variable(n->name->symbol.NameView())
+                               << " declared here";
         },
         [&](const ast::Function* n) {
-            AddNote(n->source) << "function " << style::Function << n->name->symbol.Name()
-                               << style::Plain << " declared here";
+            AddNote(n->source) << "function " << style::Function(n->name->symbol.NameView())
+                               << " declared here";
         });
 }
 
diff --git a/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc b/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
index fe16964..98aa12d 100644
--- a/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
+++ b/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
@@ -70,7 +70,7 @@
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(), R"(12:34 error: unresolved texel format 'rba8unorm'
 12:34 note: Did you mean 'rgba8unorm'?
-Possible values: 'bgra8unorm', 'r32float', 'r32sint', 'r32uint', 'rg32float', 'rg32sint', 'rg32uint', 'rgba16float', 'rgba16sint', 'rgba16uint', 'rgba32float', 'rgba32sint', 'rgba32uint', 'rgba8sint', 'rgba8snorm', 'rgba8uint', 'rgba8unorm')");
+Possible values: 'bgra8unorm', 'r32float', 'r32sint', 'r32uint', 'r8unorm', 'rg32float', 'rg32sint', 'rg32uint', 'rgba16float', 'rgba16sint', 'rgba16uint', 'rgba32float', 'rgba32sint', 'rgba32uint', 'rgba8sint', 'rgba8snorm', 'rgba8uint', 'rgba8unorm')");
 }
 
 TEST_F(ResolverUnresolvedIdentifierSuggestions, AccessMode) {
diff --git a/src/tint/lang/wgsl/resolver/validation_test.cc b/src/tint/lang/wgsl/resolver/validation_test.cc
index 809ade9..db92d22 100644
--- a/src/tint/lang/wgsl/resolver/validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/validation_test.cc
@@ -315,7 +315,7 @@
     EXPECT_FALSE(r()->Resolve());
 
     EXPECT_EQ(r()->error(),
-              R"(12:34 error: variables of type 'sampler' must not specifiy an address space)");
+              R"(12:34 error: variables of type 'sampler' must not specify an address space)");
 }
 
 TEST_F(ResolverValidationTest, AddressSpace_TextureExplicitAddressSpace) {
@@ -326,7 +326,7 @@
 
     EXPECT_EQ(
         r()->error(),
-        R"(12:34 error: variables of type 'texture_1d<f32>' must not specifiy an address space)");
+        R"(12:34 error: variables of type 'texture_1d<f32>' must not specify an address space)");
 }
 
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) {
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index 9546c27..c6412ff 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -29,6 +29,7 @@
 
 #include <algorithm>
 #include <limits>
+#include <string_view>
 #include <tuple>
 #include <utility>
 
@@ -114,22 +115,23 @@
 bool IsValidStorageTextureTexelFormat(core::TexelFormat format) {
     switch (format) {
         case core::TexelFormat::kBgra8Unorm:
-        case core::TexelFormat::kR32Uint:
-        case core::TexelFormat::kR32Sint:
         case core::TexelFormat::kR32Float:
-        case core::TexelFormat::kRg32Uint:
-        case core::TexelFormat::kRg32Sint:
+        case core::TexelFormat::kR32Sint:
+        case core::TexelFormat::kR32Uint:
+        case core::TexelFormat::kR8Unorm:
         case core::TexelFormat::kRg32Float:
-        case core::TexelFormat::kRgba8Unorm:
+        case core::TexelFormat::kRg32Sint:
+        case core::TexelFormat::kRg32Uint:
+        case core::TexelFormat::kRgba16Float:
+        case core::TexelFormat::kRgba16Sint:
+        case core::TexelFormat::kRgba16Uint:
+        case core::TexelFormat::kRgba32Float:
+        case core::TexelFormat::kRgba32Sint:
+        case core::TexelFormat::kRgba32Uint:
+        case core::TexelFormat::kRgba8Sint:
         case core::TexelFormat::kRgba8Snorm:
         case core::TexelFormat::kRgba8Uint:
-        case core::TexelFormat::kRgba8Sint:
-        case core::TexelFormat::kRgba16Uint:
-        case core::TexelFormat::kRgba16Sint:
-        case core::TexelFormat::kRgba16Float:
-        case core::TexelFormat::kRgba32Uint:
-        case core::TexelFormat::kRgba32Sint:
-        case core::TexelFormat::kRgba32Float:
+        case core::TexelFormat::kRgba8Unorm:
             return true;
         default:
             return false;
@@ -309,10 +311,9 @@
     for (auto pair : incompatible) {
         if (enabled_extensions_.Contains(pair.first) && enabled_extensions_.Contains(pair.second)) {
             AddError(source_of(pair.first))
-                << "extension " << style::Code << pair.first << style::Plain
-                << " cannot be used with extension " << style::Code << pair.second;
-            AddNote(source_of(pair.second))
-                << style::Code << pair.second << style::Plain << " enabled here";
+                << "extension " << style::Code(pair.first) << " cannot be used with extension "
+                << style::Code(pair.second);
+            AddNote(source_of(pair.second)) << style::Code(pair.second) << " enabled here";
             return false;
         }
     }
@@ -363,21 +364,18 @@
         case core::Access::kRead:
             if (!allowed_features_.features.count(
                     wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures)) {
-                AddError(source) <<
-
-                    "read-only storage textures require the "
-                    "readonly_and_readwrite_storage_textures language feature, which is not "
-                    "allowed in the current environment";
+                AddError(source) << "read-only storage textures require the "
+                                    "readonly_and_readwrite_storage_textures language feature, "
+                                    "which is not allowed in the current environment";
                 return false;
             }
             break;
         case core::Access::kReadWrite:
             if (!allowed_features_.features.count(
                     wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures)) {
-                AddError(source)
-                    << "read-write storage textures require the "
-                       "readonly_and_readwrite_storage_textures language feature, which is not "
-                       "allowed in the current environment";
+                AddError(source) << "read-write storage textures require the "
+                                    "readonly_and_readwrite_storage_textures language feature, "
+                                    "which is not allowed in the current environment";
                 return false;
             }
             break;
@@ -388,6 +386,13 @@
             return false;
     }
 
+    if (TINT_UNLIKELY(t->texel_format() == core::TexelFormat::kR8Unorm &&
+                      !enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalGraphite))) {
+        AddError(source) << style::Enum(core::TexelFormat::kR8Unorm) << " requires the "
+                         << style::Code(wgsl::Extension::kChromiumInternalGraphite) << " extension";
+        return false;
+    }
+
     if (!IsValidStorageTextureDimension(t->dim())) {
         AddError(source) << "cube dimensions for storage textures are not supported";
         return false;
@@ -429,8 +434,8 @@
                             const core::type::Type* from,
                             const Source& source) const {
     if (core::type::Type::ConversionRank(from, to) == core::type::Type::kNoConversion) {
-        AddError(source) << "cannot convert value of type " << style::Type << sem_.TypeNameOf(from)
-                         << style::Plain << " to type " << style::Type << sem_.TypeNameOf(to);
+        AddError(source) << "cannot convert value of type " << style::Type(sem_.TypeNameOf(from))
+                         << " to type " << style::Type(sem_.TypeNameOf(to));
         return false;
     }
     return true;
@@ -444,10 +449,9 @@
 
     // Value type has to match storage type
     if (storage_ty != value_type) {
-        AddError(v->source) << "cannot initialize " << style::Keyword << v->Kind() << style::Plain
-                            << " of type " << style::Type << sem_.TypeNameOf(storage_ty)
-                            << style::Plain << " with value of type " << style::Type
-                            << sem_.TypeNameOf(initializer_ty);
+        AddError(v->source) << "cannot initialize " << style::Keyword(v->Kind()) << " of type "
+                            << style::Type(sem_.TypeNameOf(storage_ty)) << " with value of type "
+                            << style::Type(sem_.TypeNameOf(initializer_ty));
         return false;
     }
 
@@ -489,17 +493,15 @@
     }
 
     auto note_usage = [&] {
-        AddNote(source) << style::Type << store_ty->FriendlyName() << style::Plain
-                        << " used in address space " << style::Enum << address_space << style::Plain
-                        << " here";
+        AddNote(source) << style::Type(store_ty->FriendlyName()) << " used in address space "
+                        << style::Enum(address_space) << " here";
     };
 
     // Among three host-shareable address spaces, f16 is supported in "uniform" and
     // "storage" address space, but not "push_constant" address space yet.
     if (Is<core::type::F16>(store_ty->DeepestElement()) &&
         address_space == core::AddressSpace::kPushConstant) {
-        AddError(source) << "using " << style::Type << "f16" << style::Plain << " in "
-                         << style::Enum << "push_constant" << style::Plain
+        AddError(source) << "using " << style::Type("f16") << " in " << style::Enum("push_constant")
                          << " address space is not implemented yet";
         return false;
     }
@@ -521,14 +523,13 @@
                 !enabled_extensions_.Contains(
                     wgsl::Extension::kChromiumInternalRelaxedUniformLayout)) {
                 AddError(m->Declaration()->source)
-                    << "the offset of a struct member of type " << style::Type
-                    << m->Type()->UnwrapRef()->FriendlyName() << style::Plain
-                    << " in address space " << style::Enum << address_space << style::Plain
-                    << " must be a multiple of " << required_align << " bytes, but "
-                    << style::Variable << member_name_of(m) << style::Plain
+                    << "the offset of a struct member of type "
+                    << style::Type(m->Type()->UnwrapRef()->FriendlyName()) << " in address space "
+                    << style::Enum(address_space) << " must be a multiple of " << required_align
+                    << " bytes, but " << style::Variable(member_name_of(m))
                     << " is currently at offset " << m->Offset() << ". Consider setting "
-                    << style::Attribute << "@align" << style::Code << "(" << required_align << ")"
-                    << style::Plain << " on this member";
+                    << style::Attribute("@align") << style::Code("(", required_align, ")")
+                    << " on this member";
 
                 AddNote(str->Declaration()->source) << "see layout of struct:\n" << str->Layout();
 
@@ -550,15 +551,14 @@
                     !enabled_extensions_.Contains(
                         wgsl::Extension::kChromiumInternalRelaxedUniformLayout)) {
                     AddError(m->Declaration()->source)
-                        << style::Enum << "uniform" << style::Plain
+                        << style::Enum("uniform")
                         << " storage requires that the number of bytes between the start of the "
                            "previous member of type struct and the current member be a "
                            "multiple of 16 bytes, but there are currently "
-                        << prev_to_curr_offset << " bytes between " << style::Variable
-                        << member_name_of(prev_member) << style::Plain << " and " << style::Variable
-                        << member_name_of(m) << style::Plain << ". Consider setting "
-                        << style::Attribute << "@align" << style::Code << "(16)" << style::Plain
-                        << " on this member";
+                        << prev_to_curr_offset << " bytes between "
+                        << style::Variable(member_name_of(prev_member)) << " and "
+                        << style::Variable(member_name_of(m)) << ". Consider setting "
+                        << style::Attribute("@align") << style::Code("(16)") << " on this member";
 
                     AddNote(str->Declaration()->source) << "see layout of struct:\n"
                                                         << str->Layout();
@@ -598,16 +598,16 @@
                            vec && vec->type()->Size() == 4) {
                     hint << "Consider using a vec4 instead.";
                 } else if (arr->ElemType()->Is<sem::Struct>()) {
-                    hint << "Consider using the " << style::Attribute << "@size" << style::Plain
+                    hint << "Consider using the " << style::Attribute("@size")
                          << " attribute on the last struct member.";
                 } else {
                     hint << "Consider wrapping the element type in a struct and using the "
-                         << style::Attribute << "@size" << style::Plain << " attribute.";
+                         << style::Attribute("@size") << " attribute.";
                 }
-                AddError(source) << style::Enum << "uniform" << style::Plain
-                                 << " storage requires that array elements are aligned to "
-                                    "16 bytes, but array element of type "
-                                 << style::Type << arr->ElemType()->FriendlyName() << style::Plain
+                AddError(source) << style::Enum("uniform")
+                                 << " storage requires that array elements are aligned to 16 "
+                                    "bytes, but array element of type "
+                                 << style::Type(arr->ElemType()->FriendlyName())
                                  << " has a stride of " << arr->Stride() << " bytes. " << hint;
                 return false;
             }
@@ -631,7 +631,7 @@
                                     ast::DisabledValidation::kIgnoreAddressSpace)) {
                 if (!local->Type()->UnwrapRef()->IsConstructible()) {
                     AddError(var->type ? var->type->source : var->source)
-                        << "function-scope " << style::Keyword << "var" << style::Plain
+                        << "function-scope " << style::Keyword("var")
                         << " must have a constructible type";
                     return false;
                 }
@@ -659,17 +659,15 @@
             if (auto* init = global->Initializer();
                 init && init->Stage() > core::EvaluationStage::kOverride) {
                 AddError(init->Declaration()->source)
-                    << "module-scope " << style::Keyword << "var" << style::Plain
-                    << " initializer must be a constant or "
-                       "override-expression";
+                    << "module-scope " << style::Keyword("var")
+                    << " initializer must be a constant or override-expression";
                 return false;
             }
 
             if (!var->declared_address_space && !global->Type()->UnwrapRef()->is_handle()) {
-                AddError(decl->source) << "module-scope " << style::Keyword << "var" << style::Plain
-                                       << " declarations that are not of texture "
-                                          "or sampler types must "
-                                          "provide an address space";
+                AddError(decl->source) << "module-scope " << style::Keyword("var")
+                                       << " declarations that are not of texture or sampler types "
+                                          "must provide an address space";
                 return false;
             }
 
@@ -684,8 +682,8 @@
     }
 
     if (global->AddressSpace() == core::AddressSpace::kFunction) {
-        AddError(decl->source) << "module-scope " << style::Keyword << "var" << style::Plain
-                               << " must not use address space " << style::Enum << "function";
+        AddError(decl->source) << "module-scope " << style::Keyword("var")
+                               << " must not use address space " << style::Enum("function");
         return false;
     }
 
@@ -697,8 +695,8 @@
             // Each resource variable must be declared with both group and binding attributes.
             if (!decl->HasBindingPoint()) {
                 AddError(decl->source)
-                    << "resource variables require " << style::Attribute << "@group" << style::Plain
-                    << " and " << style::Attribute << "@binding" << style::Plain << " attributes";
+                    << "resource variables require " << style::Attribute("@group") << " and "
+                    << style::Attribute("@binding") << " attributes";
                 return false;
             }
             break;
@@ -710,9 +708,8 @@
                 // https://gpuweb.github.io/gpuweb/wgsl/#attribute-binding
                 // Must only be applied to a resource variable
                 AddError(decl->source)
-                    << "non-resource variables must not have " << style::Attribute << "@group"
-                    << style::Plain << " or " << style::Attribute << "@binding" << style::Plain
-                    << " attributes";
+                    << "non-resource variables must not have " << style::Attribute("@group")
+                    << " or " << style::Attribute("@binding") << " attributes";
                 return false;
             }
         }
@@ -735,8 +732,8 @@
         // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
         // If the store type is a texture type or a sampler type, then the variable declaration must
         // not have a address space attribute. The address space will always be handle.
-        AddError(var->source) << "variables of type " << style::Type << sem_.TypeNameOf(store_ty)
-                              << style::Plain << " must not specifiy an address space";
+        AddError(var->source) << "variables of type " << style::Type(sem_.TypeNameOf(store_ty))
+                              << " must not specify an address space";
         return false;
     }
 
@@ -762,11 +759,10 @@
                 // Optionally has an initializer expression, if the variable is in the private or
                 // function address spaces.
                 AddError(var->source)
-                    << "var of address space " << style::Enum << v->AddressSpace() << style::Plain
-                    << " cannot have an initializer. var initializers are only supported for "
-                       "the address spaces "
-                    << style::Enum << "private" << style::Plain << " and " << style::Enum
-                    << "function";
+                    << "var of address space " << style::Enum(v->AddressSpace())
+                    << " cannot have an initializer. var initializers are only supported for the "
+                       "address spaces "
+                    << style::Enum("private") << " and " << style::Enum("function");
                 return false;
         }
     }
@@ -791,7 +787,7 @@
 
     if (!(storage_ty->IsConstructible() || storage_ty->Is<core::type::Pointer>())) {
         AddError(decl->source) << sem_.TypeNameOf(storage_ty) << " cannot be used as the type of a "
-                               << style::Keyword << "let";
+                               << style::Keyword("let");
         return false;
     }
     return true;
@@ -803,16 +799,15 @@
     auto* storage_ty = v->Type()->UnwrapRef();
 
     if (auto* init = v->Initializer(); init && init->Stage() > core::EvaluationStage::kOverride) {
-        AddError(init->Declaration()->source) << style::Keyword << "override" << style::Plain
-                                              << " initializer must be an override-expression";
+        AddError(init->Declaration()->source)
+            << style::Keyword("override") << " initializer must be an override-expression";
         return false;
     }
 
     if (auto id = v->Attributes().override_id) {
         if (auto var = override_ids.Get(*id); var && *var != v) {
             auto* attr = ast::GetAttribute<ast::IdAttribute>(v->Declaration()->attributes);
-            AddError(attr->source)
-                << style::Attribute << "@id" << style::Plain << " values must be unique";
+            AddError(attr->source) << style::Attribute("@id") << " values must be unique";
             AddNote(ast::GetAttribute<ast::IdAttribute>((*var)->Declaration()->attributes)->source)
                 << "a override with an ID of " << id->value << " was previously declared here";
             return false;
@@ -821,7 +816,7 @@
 
     if (!storage_ty->Is<core::type::Scalar>()) {
         AddError(decl->source) << sem_.TypeNameOf(storage_ty) << " cannot be used as the type of a "
-                               << style::Keyword << "override";
+                               << style::Keyword("override");
         return false;
     }
 
@@ -860,7 +855,7 @@
             }
             if (!ok) {
                 AddError(decl->source) << "function parameter of pointer type cannot be in "
-                                       << style::Enum << sc << style::Plain << " address space";
+                                       << style::Enum(sc) << " address space";
                 return false;
             }
         }
@@ -886,16 +881,14 @@
                                  ast::PipelineStage stage,
                                  const bool is_input) const {
     auto* type = storage_ty->UnwrapRef();
-    StringStream stage_name;
-    stage_name << stage;
     bool is_stage_mismatch = false;
     bool is_output = !is_input;
     auto builtin = sem_.Get(attr)->Value();
 
-    auto err_builtin_type = [&](const char* required) {
-        AddError(attr->source) << "store type of " << style::Attribute << "@builtin" << style::Code
-                               << "(" << style::Enum << builtin << style::Code << ")"
-                               << style::Plain << " must be " << style::Type << required;
+    auto err_builtin_type = [&](std::string_view required) {
+        AddError(attr->source) << "store type of " << style::Attribute("@builtin")
+                               << style::Code("(", style::Enum(builtin), ")") << " must be "
+                               << style::Type(required);
     };
 
     switch (builtin) {
@@ -989,10 +982,10 @@
         case core::BuiltinValue::kSubgroupInvocationId:
         case core::BuiltinValue::kSubgroupSize:
             if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumExperimentalSubgroups)) {
-                AddError(attr->source) << "use of " << style::Attribute << "@builtin" << style::Code
-                                       << "(" << style::Enum << builtin << style::Code << ")"
-                                       << style::Plain << " attribute requires enabling extension "
-                                       << style::Code << "chromium_experimental_subgroups";
+                AddError(attr->source) << "use of " << style::Attribute("@builtin")
+                                       << style::Code("(", style::Enum(builtin), ")")
+                                       << " attribute requires enabling extension "
+                                       << style::Code("chromium_experimental_subgroups");
                 return false;
             }
             if (!type->Is<core::type::U32>()) {
@@ -1001,8 +994,7 @@
             }
             if (stage != ast::PipelineStage::kNone && stage != ast::PipelineStage::kCompute) {
                 AddError(attr->source)
-                    << style::Attribute << "@builtin" << style::Code << "(" << style::Enum
-                    << builtin << style::Code << ")" << style::Plain
+                    << style::Attribute("@builtin") << style::Code("(", style::Enum(builtin), ")")
                     << " is only valid as a compute shader input";
                 return false;
             }
@@ -1012,9 +1004,9 @@
     }
 
     if (is_stage_mismatch) {
-        AddError(attr->source) << style::Attribute << "@builtin" << style::Code << "("
-                               << style::Enum << builtin << style::Code << ")" << style::Plain
-                               << " cannot be used for " << stage_name.str() << " shader "
+        AddError(attr->source) << style::Attribute("@builtin")
+                               << style::Code("(", style::Enum(builtin), ")")
+                               << " cannot be used for " << stage << " shader "
                                << (is_input ? "input" : "output");
         return false;
     }
@@ -1026,7 +1018,7 @@
                                      const core::type::Type* storage_ty,
                                      const ast::PipelineStage stage) const {
     if (stage == ast::PipelineStage::kCompute) {
-        AddError(attr->source) << style::Attribute << "@" << attr->Name() << style::Plain
+        AddError(attr->source) << style::Attribute("@", attr->Name())
                                << " cannot be used by compute shaders";
         return false;
     }
@@ -1039,8 +1031,8 @@
     }
 
     if (type->is_integer_scalar_or_vector() && i_type->Value() != core::InterpolationType::kFlat) {
-        AddError(attr->source) << "interpolation type must be " << style::Enum << "flat"
-                               << style::Plain << " for integral user-defined IO types";
+        AddError(attr->source) << "interpolation type must be " << style::Enum("flat")
+                               << " for integral user-defined IO types";
         return false;
     }
 
@@ -1055,7 +1047,7 @@
 bool Validator::InvariantAttribute(const ast::InvariantAttribute* attr,
                                    const ast::PipelineStage stage) const {
     if (stage == ast::PipelineStage::kCompute) {
-        AddError(attr->source) << style::Attribute << "@" << attr->Name() << style::Plain
+        AddError(attr->source) << style::Attribute("@", attr->Name())
                                << " cannot be used by compute shaders";
         return false;
     }
@@ -1070,7 +1062,7 @@
             attr,  //
             [&](const ast::WorkgroupAttribute*) {
                 if (decl->PipelineStage() != ast::PipelineStage::kCompute) {
-                    AddError(attr->source) << style::Attribute << "@workgroup_size" << style::Plain
+                    AddError(attr->source) << style::Attribute("@workgroup_size")
                                            << " is only valid for compute stages";
                     return false;
                 }
@@ -1079,7 +1071,7 @@
             [&](const ast::MustUseAttribute*) {
                 if (func->ReturnType()->Is<core::type::Void>()) {
                     AddError(attr->source)
-                        << style::Attribute << "@must_use" << style::Plain
+                        << style::Attribute("@must_use")
                         << " can only be applied to functions that return a value";
                     return false;
                 }
@@ -1115,7 +1107,7 @@
             }
         } else if (TINT_UNLIKELY(IsValidationEnabled(
                        decl->attributes, ast::DisabledValidation::kFunctionHasNoBody))) {
-            TINT_ICE() << "function " << decl->name->symbol.Name() << " has no body";
+            TINT_ICE() << "function " << decl->name->symbol.NameView() << " has no body";
         }
     }
 
@@ -1129,7 +1121,7 @@
     // a function behavior is always one of {}, or {Next}.
     if (TINT_UNLIKELY(func->Behaviors() != sem::Behaviors{} &&
                       func->Behaviors() != sem::Behavior::kNext)) {
-        auto name = decl->name->symbol.Name();
+        auto name = decl->name->symbol.NameView();
         TINT_ICE() << "function '" << name << "' behaviors are: " << func->Behaviors();
     }
 
@@ -1180,16 +1172,16 @@
                     if (pipeline_io_attribute) {
                         AddError(attr->source) << "multiple entry point IO attributes";
                         AddNote(pipeline_io_attribute->source)
-                            << "previously consumed " << style::Attribute << "@"
-                            << pipeline_io_attribute->Name();
+                            << "previously consumed "
+                            << style::Attribute("@", pipeline_io_attribute->Name());
                         return false;
                     }
                     pipeline_io_attribute = attr;
 
                     if (builtins.Contains(builtin)) {
                         AddError(decl->source)
-                            << style::Attribute << "@builtin" << style::Code << "(" << style::Enum
-                            << builtin << style::Code << ")" << style::Plain
+                            << style::Attribute("@builtin")
+                            << style::Code("(", style::Enum(builtin), ")")
                             << " appears multiple times as pipeline "
                             << (param_or_ret == ParamOrRetType::kParameter ? "input" : "output");
                         return false;
@@ -1208,8 +1200,8 @@
                     if (pipeline_io_attribute) {
                         AddError(attr->source) << "multiple entry point IO attributes";
                         AddNote(pipeline_io_attribute->source)
-                            << "previously consumed " << style::Attribute << "@"
-                            << pipeline_io_attribute->Name();
+                            << "previously consumed "
+                            << style::Attribute("@", pipeline_io_attribute->Name());
                         return false;
                     }
                     pipeline_io_attribute = attr;
@@ -1236,8 +1228,8 @@
                     if (pipeline_io_attribute) {
                         AddError(attr->source) << "multiple entry point IO attributes";
                         AddNote(pipeline_io_attribute->source)
-                            << "previously consumed " << style::Attribute << "@"
-                            << pipeline_io_attribute->Name();
+                            << "previously consumed "
+                            << style::Attribute("@", pipeline_io_attribute->Name());
                         return false;
                     }
                     pipeline_io_attribute = attr;
@@ -1285,18 +1277,18 @@
                 if (ty->is_integer_scalar_or_vector() && !interpolate_attribute) {
                     if (decl->PipelineStage() == ast::PipelineStage::kVertex &&
                         param_or_ret == ParamOrRetType::kReturnType) {
-                        AddError(source) << "integral user-defined vertex outputs must have a "
-                                         << style::Attribute << "@interpolate" << style::Code << "("
-                                         << style::Enum << "flat" << style::Code << ")"
-                                         << style::Plain << " attribute";
+                        AddError(source)
+                            << "integral user-defined vertex outputs must have a "
+                            << style::Attribute("@interpolate")
+                            << style::Code("(", style::Enum("flat"), ")") << " attribute";
                         return false;
                     }
                     if (decl->PipelineStage() == ast::PipelineStage::kFragment &&
                         param_or_ret == ParamOrRetType::kParameter) {
-                        AddError(source) << "integral user-defined fragment inputs must have a "
-                                         << style::Attribute << "@interpolate" << style::Code << "("
-                                         << style::Enum << "flat" << style::Code << ")"
-                                         << style::Plain << " attribute";
+                        AddError(source)
+                            << "integral user-defined fragment inputs must have a "
+                            << style::Attribute("@interpolate")
+                            << style::Code("(", style::Enum("flat"), ")") << " attribute";
                         return false;
                     }
                 }
@@ -1308,9 +1300,9 @@
                 // in the backend writers.
                 if (location.value_or(1) != 0) {
                     AddError(blend_src_attribute->source)
-                        << style::Attribute << "@blend_src" << style::Plain
-                        << " can only be used with " << style::Attribute << "@location"
-                        << style::Code << "(" << style::Literal << "0" << style::Code << ")";
+                        << style::Attribute("@blend_src") << " can only be used with "
+                        << style::Attribute("@location")
+                        << style::Code("(", style::Literal("0"), ")");
                     return false;
                 }
             }
@@ -1323,12 +1315,12 @@
 
             if (first_blend_src && first_location_without_blend_src) {
                 AddError(first_location_without_blend_src->source)
-                    << "use of " << style::Attribute << "@blend_src" << style::Plain
-                    << " requires all the output " << style::Attribute << "@location"
-                    << style::Plain << " attributes of the entry point to be paired with a "
-                    << style::Attribute << "@blend_src" << style::Plain << " attribute";
+                    << "use of " << style::Attribute("@blend_src") << " requires all the output "
+                    << style::Attribute("@location")
+                    << " attributes of the entry point to be paired with a "
+                    << style::Attribute("@blend_src") << " attribute";
                 AddNote(first_blend_src->source)
-                    << "use of " << style::Attribute << "@blend_src" << style::Plain << " here";
+                    << "use of " << style::Attribute("@blend_src") << " here";
                 return false;
             }
 
@@ -1338,11 +1330,10 @@
                 }
                 if (first_nonzero_location && first_blend_src) {
                     AddError(first_blend_src->source)
-                        << "pipeline cannot use both a " << style::Attribute << "@blend_src"
-                        << style::Plain << " and non-zero " << style::Attribute << "@location";
+                        << "pipeline cannot use both a " << style::Attribute("@blend_src")
+                        << " and non-zero " << style::Attribute("@location");
                     AddNote(first_nonzero_location->source)
-                        << "non-zero " << style::Attribute << "@location" << style::Plain
-                        << " declared here";
+                        << "non-zero " << style::Attribute("@location") << " declared here";
                     return false;
                 }
 
@@ -1350,21 +1341,21 @@
                                                                      blend_src.value_or(0));
                 if (!locations_and_blend_srcs.Add(location_and_blend_src)) {
                     auto& err = AddError(location_attribute->source)
-                                << style::Attribute << "@location" << style::Code << "("
-                                << style::Literal << location.value() << style::Code << ")";
+                                << style::Attribute("@location")
+                                << style::Code("(", style::Literal(location.value()), ")");
                     if (blend_src_attribute) {
-                        err << style::Attribute << " @blend_src" << style::Code << "("
-                            << style::Literal << blend_src.value() << style::Code << ")";
+                        err << style::Attribute(" @blend_src")
+                            << style::Code("(", style::Literal(blend_src.value()), ")");
                     }
-                    err << style::Plain << " appears multiple times";
+                    err << " appears multiple times";
                     return false;
                 }
             }
 
             if (color_attribute && !colors.Add(color.value())) {
                 AddError(color_attribute->source)
-                    << style::Attribute << "@color" << style::Code << "(" << style::Literal
-                    << color.value() << style::Code << ")" << style::Plain
+                    << style::Attribute("@color")
+                    << style::Code("(", style::Literal(color.value()), ")")
                     << " appears multiple times";
                 return false;
             }
@@ -1373,8 +1364,8 @@
                 if (!pipeline_io_attribute ||
                     !pipeline_io_attribute->Is<ast::LocationAttribute>()) {
                     AddError(interpolate_attribute->source)
-                        << style::Attribute << "@interpolate" << style::Plain
-                        << " can only be used with " << style::Attribute << "@location";
+                        << style::Attribute("@interpolate") << " can only be used with "
+                        << style::Attribute("@location");
                     return false;
                 }
             }
@@ -1389,9 +1380,9 @@
                 }
                 if (!has_position) {
                     AddError(invariant_attribute->source)
-                        << style::Attribute << "@invariant" << style::Plain
-                        << " must be applied to a " << style::Attribute << "@builtin" << style::Code
-                        << "(" << style::Enum << "position" << style::Code << ")";
+                        << style::Attribute("@invariant") << " must be applied to a "
+                        << style::Attribute("@builtin")
+                        << style::Code("(", style::Enum("position"), ")");
                     return false;
                 }
             }
@@ -1417,8 +1408,8 @@
                             member->Declaration()->source, param_or_ret,
                             /*is_struct_member*/ true, member->Attributes().location,
                             member->Attributes().blend_src, member->Attributes().color)) {
-                        AddNote(decl->source) << "while analyzing entry point " << style::Function
-                                              << decl->name->symbol.Name();
+                        AddNote(decl->source) << "while analyzing entry point "
+                                              << style::Function(decl->name->symbol.NameView());
                         return false;
                     }
                 }
@@ -1469,16 +1460,16 @@
             }
         }
         if (!found) {
-            AddError(decl->source) << "a vertex shader must include the " << style::Enum
-                                   << "position" << style::Plain << " builtin in its return type";
+            AddError(decl->source) << "a vertex shader must include the " << style::Enum("position")
+                                   << " builtin in its return type";
             return false;
         }
     }
 
     if (decl->PipelineStage() == ast::PipelineStage::kCompute) {
         if (!ast::HasAttribute<ast::WorkgroupAttribute>(decl->attributes)) {
-            AddError(decl->source) << "a compute shader must include " << style::Attribute
-                                   << "@workgroup_size" << style::Plain << " in its attributes";
+            AddError(decl->source) << "a compute shader must include "
+                                   << style::Attribute("@workgroup_size") << " in its attributes";
             return false;
         }
     }
@@ -1504,13 +1495,13 @@
             // Bindings must not alias within a shader stage: two different variables in the
             // resource interface of a given shader must not have the same group and binding values,
             // when considered as a pair of values.
-            auto func_name = decl->name->symbol.Name();
+            auto func_name = decl->name->symbol.NameView();
             AddError(var_decl->source)
-                << "entry point " << style::Function << func_name << style::Plain
+                << "entry point " << style::Function(func_name)
                 << " references multiple variables that use the same resource binding "
-                << style::Attribute << "@group" << style::Code << "(" << style::Literal << bp->group
-                << style::Code << ")" << style::Plain << ", " << style::Attribute << "@binding"
-                << style::Code << "(" << style::Literal << bp->binding << style::Code << ")";
+                << style::Attribute("@group") << style::Code("(", style::Literal(bp->group), ")")
+                << ", " << style::Attribute("@binding")
+                << style::Code("(", style::Literal(bp->binding), ")");
             AddNote(added.value->source) << "first resource binding usage declared here";
             return false;
         }
@@ -1547,8 +1538,8 @@
         if (auto* stmt = expr->Stmt()) {
             if (auto* decl = As<ast::VariableDeclStatement>(stmt->Declaration())) {
                 if (decl->variable->Is<ast::Const>()) {
-                    AddNote(decl->source) << "consider changing " << style::Keyword << "const"
-                                          << style::Plain << " to " << style::Keyword << "let";
+                    AddNote(decl->source) << "consider changing " << style::Keyword("const")
+                                          << " to " << style::Keyword("let");
                 }
             }
         }
@@ -1624,14 +1615,14 @@
             call->Target(),  //
             [&](const sem::Function* fn) {
                 AddError(call->Declaration()->source)
-                    << "ignoring return value of function " << style::Function
-                    << fn->Declaration()->name->symbol.Name() << style::Plain << " annotated with "
-                    << style::Attribute << "@must_use";
+                    << "ignoring return value of function "
+                    << style::Function(fn->Declaration()->name->symbol.NameView())
+                    << " annotated with " << style::Attribute("@must_use");
                 sem_.NoteDeclarationSource(fn->Declaration());
             },
             [&](const sem::BuiltinFn* b) {
                 AddError(call->Declaration()->source)
-                    << "ignoring return value of builtin " << style::Function << b->Fn();
+                    << "ignoring return value of builtin " << style::Function(b->Fn());
             },
             [&](const sem::ValueConversion*) {
                 AddError(call->Declaration()->source) << "value conversion evaluated but not used";
@@ -1745,7 +1736,7 @@
             // used instead.
             auto* builtin = call->Target()->As<sem::BuiltinFn>();
             AddError(call->Declaration()->source)
-                << "builtin function " << style::Function << builtin->Fn() << style::Plain
+                << "builtin function " << style::Function(builtin->Fn())
                 << " does not return a value";
             return false;
         }
@@ -1851,8 +1842,8 @@
     if (extension != wgsl::Extension::kUndefined) {
         if (!enabled_extensions_.Contains(extension)) {
             AddError(call->Declaration()->source)
-                << "cannot call built-in function " << style::Function << builtin->Fn()
-                << style::Plain << " without extension " << extension;
+                << "cannot call built-in function " << style::Function(builtin->Fn())
+                << " without extension " << extension;
             return false;
         }
     }
@@ -1861,8 +1852,8 @@
     if (feature != wgsl::LanguageFeature::kUndefined) {
         if (!allowed_features_.features.count(feature)) {
             AddError(call->Declaration()->source)
-                << "built-in function " << style::Function << builtin->Fn() << style::Plain
-                << " requires the " << style::Code << wgsl::ToString(feature) << style::Plain
+                << "built-in function " << style::Function(builtin->Fn()) << " requires the "
+                << style::Code(wgsl::ToString(feature))
                 << " language feature, which is not allowed in the current environment";
             return false;
         }
@@ -1874,8 +1865,8 @@
 bool Validator::CheckF16Enabled(const Source& source) const {
     // Validate if f16 type is allowed.
     if (!enabled_extensions_.Contains(wgsl::Extension::kF16)) {
-        AddError(source) << style::Type << "f16" << style::Plain << " type used without "
-                         << style::Code << "f16" << style::Plain << " extension enabled";
+        AddError(source) << style::Type("f16") << " type used without " << style::Code("f16")
+                         << " extension enabled";
         return false;
     }
     return true;
@@ -1904,7 +1895,7 @@
         AddError(decl->source) << "too "
                                << (more ? std::string("many") : std::string("few")) +
                                       " arguments in call to "
-                               << style::Function << name << style::Plain << ", expected "
+                               << style::Function(name) << ", expected "
                                << target->Parameters().Length() << ", got "
                                << call->Arguments().Length();
         return false;
@@ -1918,9 +1909,9 @@
 
         if (param_type != arg_type) {
             AddError(arg_expr->source) << "type mismatch for argument " << (i + 1) << " in call to "
-                                       << style::Function << name << style::Plain << ", expected "
-                                       << style::Type << sem_.TypeNameOf(param_type) << style::Plain
-                                       << ", got " << style::Type << sem_.TypeNameOf(arg_type);
+                                       << style::Function(name) << ", expected "
+                                       << style::Type(sem_.TypeNameOf(param_type)) << ", got "
+                                       << style::Type(sem_.TypeNameOf(arg_type));
             return false;
         }
 
@@ -1964,8 +1955,8 @@
             // https://gpuweb.github.io/gpuweb/wgsl/#function-call-expr
             // If the called function does not return a value, a function call
             // statement should be used instead.
-            AddError(decl->source) << "function " << style::Function << name << style::Plain
-                                   << " does not return a value";
+            AddError(decl->source)
+                << "function " << style::Function(name) << " does not return a value";
             return false;
         }
     }
@@ -1994,8 +1985,8 @@
             if (member->Type() != value_ty->UnwrapRef()) {
                 AddError(value->source)
                     << "type in structure constructor does not match struct member type: expected "
-                    << style::Type << sem_.TypeNameOf(member->Type()) << style::Plain << ", found "
-                    << style::Type << sem_.TypeNameOf(value_ty);
+                    << style::Type(sem_.TypeNameOf(member->Type())) << ", found "
+                    << style::Type(sem_.TypeNameOf(value_ty));
                 return false;
             }
         }
@@ -2011,9 +2002,9 @@
         auto* value_ty = sem_.TypeOf(value)->UnwrapRef();
         if (core::type::Type::ConversionRank(value_ty, elem_ty) ==
             core::type::Type::kNoConversion) {
-            AddError(value->source) << style::Type << sem_.TypeNameOf(value_ty) << style::Plain
-                                    << " cannot be used to construct an array of " << style::Type
-                                    << sem_.TypeNameOf(elem_ty);
+            AddError(value->source) << style::Type(sem_.TypeNameOf(value_ty))
+                                    << " cannot be used to construct an array of "
+                                    << style::Type(sem_.TypeNameOf(elem_ty));
             return false;
         }
     }
@@ -2051,10 +2042,9 @@
 
 bool Validator::Vector(const core::type::Type* el_ty, const Source& source) const {
     if (!el_ty->Is<core::type::Scalar>()) {
-        AddError(source) << "vector element type must be " << style::Type << "bool" << style::Plain
-                         << ", " << style::Type << "f32" << style::Plain << ", " << style::Type
-                         << "f16" << style::Plain << ", " << style::Type << "i32" << style::Plain
-                         << " or " << style::Type << "u32";
+        AddError(source) << "vector element type must be " << style::Type("bool") << ", "
+                         << style::Type("f32") << ", " << style::Type("f16") << ", "
+                         << style::Type("i32") << " or " << style::Type("u32");
         return false;
     }
     return true;
@@ -2062,8 +2052,8 @@
 
 bool Validator::Matrix(const core::type::Type* el_ty, const Source& source) const {
     if (!el_ty->is_float_scalar()) {
-        AddError(source) << "matrix element type must be " << style::Type << "f32" << style::Plain
-                         << " or " << style::Type << "f16";
+        AddError(source) << "matrix element type must be " << style::Type("f32") << " or "
+                         << style::Type("f16");
         return false;
     }
     return true;
@@ -2073,12 +2063,13 @@
     auto backtrace = [&](const sem::Function* func, const sem::Function* entry_point) {
         if (func != entry_point) {
             TraverseCallChain(entry_point, func, [&](const sem::Function* f) {
-                AddNote(f->Declaration()->source) << "called by function " << style::Function
-                                                  << f->Declaration()->name->symbol.Name();
+                AddNote(f->Declaration()->source)
+                    << "called by function "
+                    << style::Function(f->Declaration()->name->symbol.NameView());
             });
             AddNote(entry_point->Declaration()->source)
-                << "called by entry point " << style::Function
-                << entry_point->Declaration()->name->symbol.Name();
+                << "called by entry point "
+                << style::Function(entry_point->Declaration()->name->symbol.NameView());
         }
     };
 
@@ -2091,7 +2082,7 @@
                     break;
                 }
             }
-            AddError(source) << "var with " << style::Enum << var->AddressSpace() << style::Plain
+            AddError(source) << "var with " << style::Enum(var->AddressSpace())
                              << " address space cannot be used by " << stage << " pipeline stage";
             AddNote(var->Declaration()->source) << "variable is declared here";
             backtrace(func, entry_point);
@@ -2299,7 +2290,7 @@
                 [&](const ast::StructMemberSizeAttribute*) {
                     if (!member->Type()->HasCreationFixedFootprint()) {
                         AddError(attr->source)
-                            << style::Attribute << "@size" << style::Plain
+                            << style::Attribute("@size")
                             << " can only be applied to members where the member's type size can "
                                "be fully determined at shader creation time";
                         return false;
@@ -2314,8 +2305,7 @@
 
         if (invariant_attribute && !has_position) {
             AddError(invariant_attribute->source)
-                << style::Attribute << "@invariant" << style::Plain
-                << " must be applied to a position builtin";
+                << style::Attribute("@invariant") << " must be applied to a position builtin";
             return false;
         }
 
@@ -2325,17 +2315,16 @@
             // backend writers.
             if (member->Attributes().location.value_or(1) != 0) {
                 AddError(blend_src_attribute->source)
-                    << style::Attribute << "@blend_src" << style::Plain << " can only be used with "
-                    << style::Attribute << "@location" << style::Code << "(" << style::Literal
-                    << "0" << style::Code << ")";
+                    << style::Attribute("@blend_src") << " can only be used with "
+                    << style::Attribute("@location") << style::Code("(", style::Literal("0"), ")");
                 return false;
             }
         }
 
         if (interpolate_attribute && !location_attribute) {
             AddError(interpolate_attribute->source)
-                << style::Attribute << "@interpolate" << style::Plain << " can only be used with "
-                << style::Attribute << "@location";
+                << style::Attribute("@interpolate") << " can only be used with "
+                << style::Attribute("@location");
             return false;
         }
 
@@ -2346,13 +2335,13 @@
 
             if (!locations_and_blend_srcs.Add(std::make_pair(location, blend_src))) {
                 auto& err = AddError(location_attribute->source)
-                            << style::Attribute << "@location" << style::Code << "("
-                            << style::Literal << location << style::Code << ")";
+                            << style::Attribute("@location")
+                            << style::Code("(", style::Literal(location), ")");
                 if (blend_src) {
-                    err << style::Attribute << " @blend_src" << style::Code << "(" << style::Literal
-                        << blend_src.value() << style::Code << ")";
+                    err << style::Attribute(" @blend_src")
+                        << style::Code("(", style::Literal(blend_src.value()), ")");
                 }
-                err << style::Plain << " appears multiple times";
+                err << " appears multiple times";
                 return false;
             }
         }
@@ -2361,8 +2350,8 @@
             uint32_t color = member->Attributes().color.value();
             if (!colors.Add(color)) {
                 AddError(color_attribute->source)
-                    << style::Attribute << "@color" << style::Code << "(" << style::Literal << color
-                    << style::Code << ")" << style::Plain << " appears multiple times";
+                    << style::Attribute("@color") << style::Code("(", style::Literal(color), ")")
+                    << " appears multiple times";
                 return false;
             }
         }
@@ -2376,17 +2365,17 @@
                                   ast::PipelineStage stage,
                                   const Source& source) const {
     if (stage == ast::PipelineStage::kCompute) {
-        AddError(attr->source) << style::Attribute << "@" << attr->Name() << style::Plain
+        AddError(attr->source) << style::Attribute("@", attr->Name())
                                << " cannot be used by compute shaders";
         return false;
     }
 
     if (!type->is_numeric_scalar_or_vector()) {
         std::string invalid_type = sem_.TypeNameOf(type);
-        AddError(source) << "cannot apply " << style::Attribute << "@location" << style::Plain
-                         << " to declaration of type " << style::Type << invalid_type;
+        AddError(source) << "cannot apply " << style::Attribute("@location")
+                         << " to declaration of type " << style::Type(invalid_type);
         AddNote(attr->source)
-            << style::Attribute << "@location" << style::Plain
+            << style::Attribute("@location")
             << " must only be applied to declarations of numeric scalar or numeric vector type";
         return false;
     }
@@ -2400,9 +2389,9 @@
                                const Source& source,
                                const std::optional<bool> is_input) const {
     if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumExperimentalFramebufferFetch)) {
-        AddError(attr->source) << "use of " << style::Attribute << "@color" << style::Plain
-                               << " requires enabling extension " << style::Code
-                               << "chromium_experimental_framebuffer_fetch";
+        AddError(attr->source) << "use of " << style::Attribute("@color")
+                               << " requires enabling extension "
+                               << style::Code("chromium_experimental_framebuffer_fetch");
         return false;
     }
 
@@ -2410,17 +2399,17 @@
         stage != ast::PipelineStage::kNone && stage != ast::PipelineStage::kFragment;
     bool is_output = !is_input.value_or(true);
     if (is_stage_non_fragment || is_output) {
-        AddError(attr->source) << style::Attribute << "@color" << style::Plain
+        AddError(attr->source) << style::Attribute("@color")
                                << " can only be used for fragment shader input";
         return false;
     }
 
     if (!type->is_numeric_scalar_or_vector()) {
         std::string invalid_type = sem_.TypeNameOf(type);
-        AddError(source) << "cannot apply " << style::Attribute << "@color" << style::Plain
-                         << " to declaration of type " << style::Type << invalid_type;
+        AddError(source) << "cannot apply " << style::Attribute("@color")
+                         << " to declaration of type " << style::Type(invalid_type);
         AddNote(attr->source)
-            << style::Attribute << "@color" << style::Plain
+            << style::Attribute("@color")
             << " must only be applied to declarations of numeric scalar or numeric vector type";
         return false;
     }
@@ -2432,9 +2421,9 @@
                                   ast::PipelineStage stage,
                                   const std::optional<bool> is_input) const {
     if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalDualSourceBlending)) {
-        AddError(attr->source) << "use of " << style::Attribute << "@blend_src" << style::Plain
-                               << " requires enabling extension " << style::Code
-                               << "chromium_internal_dual_source_blending";
+        AddError(attr->source) << "use of " << style::Attribute("@blend_src")
+                               << " requires enabling extension "
+                               << style::Code("chromium_internal_dual_source_blending");
         return false;
     }
 
@@ -2442,7 +2431,7 @@
         stage != ast::PipelineStage::kNone && stage != ast::PipelineStage::kFragment;
     bool is_output = is_input.value_or(false);
     if (is_stage_non_fragment || is_output) {
-        AddError(attr->source) << style::Attribute << "@" << attr->Name() << style::Plain
+        AddError(attr->source) << style::Attribute("@", attr->Name())
                                << " can only be used for fragment shader output";
         return false;
     }
@@ -2456,9 +2445,9 @@
                        sem::Statement* current_statement) const {
     if (func_type->UnwrapRef() != ret_type) {
         AddError(ret->source)
-            << "return statement type must match its function return type, returned " << style::Type
-            << sem_.TypeNameOf(ret_type) << style::Plain << ", expected " << style::Type
-            << sem_.TypeNameOf(func_type);
+            << "return statement type must match its function return type, returned "
+            << style::Type(sem_.TypeNameOf(ret_type)) << ", expected "
+            << style::Type(sem_.TypeNameOf(func_type));
         return false;
     }
 
@@ -2561,9 +2550,8 @@
             !ty->IsAnyOf<core::type::Pointer, core::type::Texture, core::type::Sampler,
                          core::type::AbstractNumeric>()) {
             AddError(rhs->source)
-                << "cannot assign " << style::Type << sem_.TypeNameOf(rhs_ty) << style::Plain
-                << " to " << style::Code << "_" << style::Plain << ". " << style::Code << "_"
-                << style::Plain
+                << "cannot assign " << style::Type(sem_.TypeNameOf(rhs_ty)) << " to "
+                << style::Code("_") << ". " << style::Code("_")
                 << " can only be assigned a constructible, pointer, texture or sampler type";
             return false;
         }
@@ -2590,20 +2578,17 @@
                             user->Variable()->Declaration(),  //
                             [&](const ast::Let* v) {
                                 AddNote(user->Declaration()->source)
-                                    << style::Variable << "let" << style::Plain
-                                    << " variables are immutable";
+                                    << style::Variable("let") << " variables are immutable";
                                 sem_.NoteDeclarationSource(v);
                             },
                             [&](const ast::Const* v) {
                                 AddNote(user->Declaration()->source)
-                                    << style::Variable << "const" << style::Plain
-                                    << " variables are immutable";
+                                    << style::Variable("const") << " variables are immutable";
                                 sem_.NoteDeclarationSource(v);
                             },
                             [&](const ast::Override* v) {
                                 AddNote(user->Declaration()->source)
-                                    << style::Variable << "override" << style::Plain
-                                    << " variables are immutable";
+                                    << style::Variable("override") << " variables are immutable";
                                 sem_.NoteDeclarationSource(v);
                             },
                             [&](const ast::Parameter* v) {
@@ -2623,8 +2608,8 @@
 
     // Value type has to match storage type
     if (storage_ty != value_type) {
-        AddError(a->source) << "cannot assign " << style::Type << sem_.TypeNameOf(rhs_ty)
-                            << style::Plain << " to " << style::Type << sem_.TypeNameOf(lhs_ty);
+        AddError(a->source) << "cannot assign " << style::Type(sem_.TypeNameOf(rhs_ty)) << " to "
+                            << style::Type(sem_.TypeNameOf(lhs_ty));
         return false;
     }
     if (!storage_ty->IsConstructible()) {
@@ -2632,8 +2617,8 @@
         return false;
     }
     if (lhs_ref->Access() == core::Access::kRead) {
-        AddError(a->source) << "cannot store into a read-only type " << style::Type
-                            << sem_.RawTypeNameOf(lhs_ty);
+        AddError(a->source) << "cannot store into a read-only type "
+                            << style::Type(sem_.RawTypeNameOf(lhs_ty));
         return false;
     }
     return true;
@@ -2653,11 +2638,11 @@
                 return true;
             },
             [&](const ast::Let*) {
-                AddError(lhs->source) << "cannot modify " << style::Keyword << "let";
+                AddError(lhs->source) << "cannot modify " << style::Keyword("let");
                 return true;
             },
             [&](const ast::Override*) {
-                AddError(lhs->source) << "cannot modify " << style::Keyword << "override";
+                AddError(lhs->source) << "cannot modify " << style::Keyword("override");
                 return true;
             });
         if (errored) {
@@ -2670,8 +2655,8 @@
     auto* lhs_ref = lhs_ty->As<core::type::Reference>();
     if (!lhs_ref) {
         // LHS is not a reference, so it has no storage.
-        AddError(lhs->source) << "cannot modify value of type " << style::Type
-                              << sem_.TypeNameOf(lhs_ty);
+        AddError(lhs->source) << "cannot modify value of type "
+                              << style::Type(sem_.TypeNameOf(lhs_ty));
         return false;
     }
 
@@ -2682,8 +2667,8 @@
     }
 
     if (lhs_ref->Access() == core::Access::kRead) {
-        AddError(inc->source) << "cannot modify read-only type " << style::Type
-                              << sem_.RawTypeNameOf(lhs_ty);
+        AddError(inc->source) << "cannot modify read-only type "
+                              << style::Type(sem_.RawTypeNameOf(lhs_ty));
         return false;
     }
     return true;
@@ -2721,8 +2706,8 @@
         if (!diag_added && diag_added.value->severity != dc->severity) {
             AddError(dc->rule_name->source) << "conflicting diagnostic " << use;
             AddNote(diag_added.value->rule_name->source)
-                << "severity of " << style::Code << dc->rule_name->String() << style::Plain
-                << " set to " << style::Code << dc->severity << style::Plain << " here";
+                << "severity of " << style::Code(dc->rule_name->String()) << " set to "
+                << style::Code(dc->severity) << " here";
             return false;
         }
     }
@@ -2756,11 +2741,9 @@
 }
 
 void Validator::RaiseArrayWithOverrideCountError(const Source& source) const {
-    AddError(source) << style::Type << "array" << style::Plain << " with an " << style::Keyword
-                     << "override" << style::Plain
-                     << " element count can only be used as the store type of a " << style::Keyword
-                     << "var" << style::Code << "<" << style::Enum << "workgroup" << style::Code
-                     << ">";
+    AddError(source) << style::Type("array") << " with an " << style::Keyword("override")
+                     << " element count can only be used as the store type of a "
+                     << style::Keyword("var") << style::Code("<", style::Enum("workgroup"), ">");
 }
 
 std::string Validator::VectorPretty(uint32_t size, const core::type::Type* element_type) const {
@@ -2784,20 +2767,18 @@
                     using Allowed = std::tuple<core::type::I32, core::type::U32, core::type::F32>;
                     if (TINT_UNLIKELY(!member->Type()->TypeInfo().IsAnyOfTuple<Allowed>())) {
                         AddError(member->Declaration()->source)
-                            << style::Keyword << "struct" << style::Plain << " members used in the "
-                            << style::Enum << "pixel_local" << style::Plain
-                            << " address space can only be of the type " << style::Type << "i32"
-                            << style::Plain << ", " << style::Type << "u32" << style::Plain
-                            << " or " << style::Type << "f32";
-                        AddNote(source)
-                            << style::Keyword << "struct " << style::Type << str->Name().Name()
-                            << style::Plain << " used in the " << style::Enum << "pixel_local"
-                            << style::Plain << " address space here";
+                            << style::Keyword("struct") << " members used in the "
+                            << style::Enum("pixel_local")
+                            << " address space can only be of the type " << style::Type("i32")
+                            << ", " << style::Type("u32") << " or " << style::Type("f32");
+                        AddNote(source) << style::Keyword("struct ")
+                                        << style::Type(str->Name().Name()) << " used in the "
+                                        << style::Enum("pixel_local") << " address space here";
                         return false;
                     }
                 }
             } else if (TINT_UNLIKELY(!store_ty->TypeInfo().Is<core::type::Struct>())) {
-                AddError(source) << style::Enum << "pixel_local" << style::Plain
+                AddError(source) << style::Enum("pixel_local")
                                  << " variable only support struct storage types";
                 return false;
             }
@@ -2807,19 +2788,18 @@
                                   wgsl::Extension::kChromiumExperimentalPushConstant) &&
                               IsValidationEnabled(attributes,
                                                   ast::DisabledValidation::kIgnoreAddressSpace))) {
-                AddError(source) << "use of variable address space " << style::Enum
-                                 << "push_constant" << style::Plain
-                                 << " requires enabling extension " << style::Code
-                                 << "chromium_experimental_push_constant";
+                AddError(source) << "use of variable address space " << style::Enum("push_constant")
+                                 << " requires enabling extension "
+                                 << style::Code("chromium_experimental_push_constant");
                 return false;
             }
             break;
         case core::AddressSpace::kStorage:
             if (TINT_UNLIKELY(access == core::Access::kWrite)) {
                 // The access mode for the storage address space can only be 'read' or 'read_write'.
-                AddError(source) << "access mode " << style::Enum << "write" << style::Plain
-                                 << " is not valid for the " << style::Enum << "storage"
-                                 << style::Plain << " address space";
+                AddError(source) << "access mode " << style::Enum("write")
+                                 << " is not valid for the " << style::Enum("storage")
+                                 << " address space";
                 return false;
             }
             break;
@@ -2831,15 +2811,15 @@
         StyledText err;
         if (address_space != core::AddressSpace::kStorage &&
             address_space != core::AddressSpace::kWorkgroup) {
-            AddError(source) << style::Type << "atomic" << style::Plain << " variables must have "
-                             << style::Enum << "storage" << style::Plain << " or " << style::Enum
-                             << "workgroup" << style::Plain << " address space";
+            AddError(source) << style::Type("atomic") << " variables must have "
+                             << style::Enum("storage") << " or " << style::Enum("workgroup")
+                             << " address space";
             return true;
         }
         if (address_space == core::AddressSpace::kStorage && access != core::Access::kReadWrite) {
-            AddError(source) << "atomic variables in " << style::Enum << "storage" << style::Plain
-                             << " address space must have " << style::Enum << "read_write"
-                             << style::Plain << " access mode";
+            AddError(source) << "atomic variables in " << style::Enum("storage")
+                             << " address space must have " << style::Enum("read_write")
+                             << " access mode";
             return true;
         }
         return false;
@@ -2849,8 +2829,8 @@
         if (auto atomic_use = atomic_composite_info_.Get(store_ty)) {
             if (TINT_UNLIKELY(atomic_error())) {
                 AddNote(**atomic_use)
-                    << "atomic sub-type of " << style::Type << sem_.TypeNameOf(store_ty)
-                    << style::Plain << " is declared here";
+                    << "atomic sub-type of " << style::Type(sem_.TypeNameOf(store_ty))
+                    << " is declared here";
                 return false;
             }
         }
@@ -2890,29 +2870,31 @@
             }
 
             AddError(ep->Declaration()->source)
-                << "entry point " << style::Function << ep->Declaration()->name->symbol.Name()
-                << style::Plain << " uses two different " << style::Enum << space << style::Plain
-                << " variables.";
-            AddNote(var->Declaration()->source) << "first " << style::Enum << space << style::Plain
-                                                << " variable declaration is here";
+                << "entry point " << style::Function(ep->Declaration()->name->symbol.NameView())
+                << " uses two different " << style::Enum(space) << " variables.";
+            AddNote(var->Declaration()->source)
+                << "first " << style::Enum(space) << " variable declaration is here";
             if (func != ep) {
                 TraverseCallChain(ep, func, [&](const sem::Function* f) {
-                    AddNote(f->Declaration()->source) << "called by function " << style::Function
-                                                      << f->Declaration()->name->symbol.Name();
+                    AddNote(f->Declaration()->source)
+                        << "called by function "
+                        << style::Function(f->Declaration()->name->symbol.NameView());
                 });
-                AddNote(ep->Declaration()->source) << "called by entry point " << style::Function
-                                                   << ep->Declaration()->name->symbol.Name();
+                AddNote(ep->Declaration()->source)
+                    << "called by entry point "
+                    << style::Function(ep->Declaration()->name->symbol.NameView());
             }
             AddNote(seen_var->Declaration()->source)
-                << "second " << style::Enum << space << style::Plain
-                << " variable declaration is here";
+                << "second " << style::Enum(space) << " variable declaration is here";
             if (seen_func != ep) {
                 TraverseCallChain(ep, seen_func, [&](const sem::Function* f) {
-                    AddNote(f->Declaration()->source) << "called by function " << style::Function
-                                                      << f->Declaration()->name->symbol.Name();
+                    AddNote(f->Declaration()->source)
+                        << "called by function "
+                        << style::Function(f->Declaration()->name->symbol.NameView());
                 });
-                AddNote(ep->Declaration()->source) << "called by entry point " << style::Function
-                                                   << ep->Declaration()->name->symbol.Name();
+                AddNote(ep->Declaration()->source)
+                    << "called by entry point "
+                    << style::Function(ep->Declaration()->name->symbol.NameView());
             }
             return false;
         }
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index b851cdb..e61ab87 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -197,7 +197,8 @@
 
 // https://gpuweb.github.io/gpuweb/wgsl/#texel-formats
 match f32_texel_format
-  : texel_format.bgra8unorm
+  : texel_format.r8unorm
+  | texel_format.bgra8unorm
   | texel_format.rgba8unorm
   | texel_format.rgba8snorm
   | texel_format.rgba16float
diff --git a/src/tint/lang/wgsl/writer/raise/rename_conflicts.cc b/src/tint/lang/wgsl/writer/raise/rename_conflicts.cc
index a6072ed..c86f3b1 100644
--- a/src/tint/lang/wgsl/writer/raise/rename_conflicts.cc
+++ b/src/tint/lang/wgsl/writer/raise/rename_conflicts.cc
@@ -68,6 +68,11 @@
 
         RegisterModuleScopeDecls();
 
+        // Process the types
+        for (auto* ty : ir.Types()) {
+            EnsureResolvable(ty);
+        }
+
         // Process the module-scope variable declarations
         for (auto* inst : *ir.root_block) {
             Process(inst);
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 8b28870..b2823db 100644
--- a/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc
+++ b/src/tint/lang/wgsl/writer/raise/rename_conflicts_test.cc
@@ -1101,5 +1101,55 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(IRToProgramRenameConflictsTest, Conflict_StructMember_ShadowedBy_Fn) {
+    auto* s = ty.Struct(b.ir.symbols.New("s"), {{b.ir.symbols.New("f"), ty.f32()}});
+
+    b.Append(mod.root_block, [&] {  //
+        b.Var(ty.ptr<private_>(s));
+    });
+
+    auto* fn = b.Function("f32", ty.void_());
+    b.Append(fn->Block(), [&] {  //
+        b.Return(fn);
+    });
+
+    auto* src = R"(
+s = struct @align(4) {
+  f:f32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<private, s, read_write> = var
+}
+
+%f32 = func():void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+s = struct @align(4) {
+  f:f32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<private, s, read_write> = var
+}
+
+%f32_1 = func():void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)";
+
+    Run();
+
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::wgsl::writer::raise
diff --git a/src/tint/utils/diagnostic/formatter.cc b/src/tint/utils/diagnostic/formatter.cc
index 664c4f8..9fd279b 100644
--- a/src/tint/utils/diagnostic/formatter.cc
+++ b/src/tint/utils/diagnostic/formatter.cc
@@ -143,12 +143,12 @@
     }
 
     if (!prefix.IsEmpty()) {
-        text << style::Plain << ": ";
+        text << style::Plain(": ");
     }
     text << style::Bold << diag.message;
 
     if (style_.print_line && src.file && rng.begin.line > 0) {
-        text << style::Plain << "\n";
+        text << style::Plain("\n");
 
         for (size_t line_num = rng.begin.line;
              (line_num <= rng.end.line) && (line_num <= src.file->content.lines.size());
@@ -168,7 +168,7 @@
                 }
             }
 
-            text << style::Plain << "\n";
+            text << style::Plain("\n");
 
             // If the line contains non-ascii characters, then we cannot assume that
             // a single utf8 code unit represents a single glyph, so don't attempt to
@@ -206,7 +206,7 @@
                 // Middle of multi-line
                 text.Repeat('^', num_glyphs(1, line_len + 1));
             }
-            text << style::Plain << "\n";
+            text << style::Plain("\n");
         }
     }
 }
diff --git a/src/tint/utils/templates/intrinsic_table_data.tmpl.inc b/src/tint/utils/templates/intrinsic_table_data.tmpl.inc
index 0308082..9ea5a45 100644
--- a/src/tint/utils/templates/intrinsic_table_data.tmpl.inc
+++ b/src/tint/utils/templates/intrinsic_table_data.tmpl.inc
@@ -291,11 +291,11 @@
 {{- end  }}
 
 {{- if .DisplayName }}
-    out {{range SplitDisplayName .DisplayName}} << style::Type << {{.}}{{end}};
+    out << style::Type({{range $i, $n := SplitDisplayName .DisplayName}}{{if $i}}, {{end}}{{$n}}{{end}});
 {{- else if .TemplateParams }}
-    out << style::Type << "{{.Name}}" << style::Code << "<"{{template "AppendTemplateParamNames" .TemplateParams}} << style::Code << ">";
+    out << style::Type("{{.Name}}", "<", {{template "AppendTemplateParamNames" .TemplateParams}}, ">");
 {{- else }}
-    out << style::Type << "{{.Name}}";
+    out << style::Type("{{.Name}}");
 {{- end  }}
   }
 };
@@ -322,8 +322,8 @@
     // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
 {{  range .Types -}}
 {{-   if      IsFirstIn . $.Types }} k{{PascalCase .Name}}Matcher.print(nullptr, out);
-{{-   else if IsLastIn  . $.Types }} out << TextStyle{} << " or "; k{{PascalCase .Name}}Matcher.print(nullptr, out);
-{{-   else                        }} out << TextStyle{} << ", "; k{{PascalCase .Name}}Matcher.print(nullptr, out);
+{{-   else if IsLastIn  . $.Types }} out << style::Plain(" or "); k{{PascalCase .Name}}Matcher.print(nullptr, out);
+{{-   else                        }} out << style::Plain(", "); k{{PascalCase .Name}}Matcher.print(nullptr, out);
 {{-   end -}}
 {{- end -}}
   }
@@ -362,9 +362,9 @@
 /* print */ [](MatchState*, StyledText& out) {
   out
 {{- range .Options -}}
-{{-   if      IsFirstIn . $.Options }}<< style::Enum << "{{.Name}}"
-{{-   else if IsLastIn  . $.Options }}<< TextStyle{} << " or " << style::Enum << "{{.Name}}"
-{{-   else                          }}<< TextStyle{} << ", " << style::Enum << "{{.Name}}"
+{{-   if      IsFirstIn . $.Options }}<< style::Enum("{{.Name}}")
+{{-   else if IsLastIn  . $.Options }}<< style::Plain(" or ") << style::Enum("{{.Name}}")
+{{-   else                          }}<< style::Plain(", ") << style::Enum("{{.Name}}")
 {{-   end -}}
 {{- end -}};
   }
@@ -479,8 +479,8 @@
 {{-                      define "AppendTemplateParamNames"                   -}}
 {{- /* ------------------------------------------------------------------ */ -}}
 {{-   range $i, $ := . -}}
-{{-     if $i }} << style::Code << ", " << style::Type << {{.Name}}
-{{-     else  }} << style::Type << {{.Name}}
+{{-     if $i }}, ", ", {{.Name}}
+{{-     else  }}{{.Name}}
 {{-     end  -}}
 {{-   end -}}
 {{- end -}}
diff --git a/src/tint/utils/text/styled_text.cc b/src/tint/utils/text/styled_text.cc
index 448ecdf..76d33cf 100644
--- a/src/tint/utils/text/styled_text.cc
+++ b/src/tint/utils/text/styled_text.cc
@@ -58,12 +58,14 @@
 }
 
 StyledText& StyledText::SetStyle(TextStyle style) {
-    if (spans_.Back().style != style) {
-        if (spans_.Back().length == 0) {
-            spans_.Back().style = style;
-        } else {
-            spans_.Push(Span{style});
-        }
+    if (spans_.Back().style == style) {
+        return *this;
+    }
+    if (spans_.Back().length == 0) {
+        spans_.Pop();
+    }
+    if (spans_.IsEmpty() || spans_.Back().style != style) {
+        spans_.Push(Span{style});
     }
     return *this;
 }
diff --git a/src/tint/utils/text/styled_text.h b/src/tint/utils/text/styled_text.h
index f46ced1..21a82aa 100644
--- a/src/tint/utils/text/styled_text.h
+++ b/src/tint/utils/text/styled_text.h
@@ -30,6 +30,8 @@
 
 #include <string>
 #include <string_view>
+#include <tuple>
+#include <type_traits>
 #include <utility>
 
 #include "src/tint/utils/containers/vector.h"
@@ -70,6 +72,9 @@
     /// Clears the text and restore the default style
     void Clear();
 
+    /// @returns the TextStyle of all future writes to this StyledText
+    TextStyle Style() const { return spans_.Back().style; }
+
     /// Sets the style for all future writes to this StyledText
     StyledText& SetStyle(TextStyle style);
 
@@ -77,6 +82,8 @@
     std::string Plain() const;
 
     /// Appends the styled text of @p other to this StyledText.
+    /// @note: Unlike `operator<<(const StyledText&)`, this StyledText's previous style will *not*
+    /// be automatically restored.
     void Append(const StyledText& other);
 
     /// repeat queues the character @p c to be written to the StyledText n times.
@@ -85,15 +92,30 @@
     /// @returns this StyledText so calls can be chained.
     StyledText& Repeat(char c, size_t n);
 
-    /// operator<<() appends @p value to the StyledText.
-    /// @p value can be a StyledText to change the style of future appends.
+    /// operator<<() appends @p value to this StyledText.
+    /// @p value which is one of:
+    ///  * a TextStyle, which sets the style of all future appends. Note that this fully replaces
+    ///    the old style (no style combinations / overlays).
+    ///  * a ScopedTextStyle, which will will use the span's style for each value, and then this
+    ///    StyledText's previous style will be restored.
+    ///  * a StyledText, which will be appended to this StyledText, and then this StyledText's
+    ///    previous style will be restored.
+    ///  * any other type will be stringified and appended to this StyledText using the current
+    ///    TextStyle.
     template <typename VALUE>
     StyledText& operator<<(VALUE&& value) {
         using T = std::decay_t<VALUE>;
         if constexpr (std::is_same_v<T, TextStyle>) {
             SetStyle(std::forward<VALUE>(value));
         } else if constexpr (std::is_same_v<T, StyledText>) {
+            auto old_style = Style();
             Append(value);
+            *this << old_style;
+        } else if constexpr (IsScopedTextStyle<T>) {
+            auto old_style = Style();
+            std::apply([&](auto&&... values) { ((*this << value.style << values), ...); },
+                       value.values);
+            *this << old_style;
         } else {
             uint32_t offset = stream_.tellp();
             stream_ << value;
diff --git a/src/tint/utils/text/text_style.h b/src/tint/utils/text/text_style.h
index 84deef9..a9ce9e1 100644
--- a/src/tint/utils/text/text_style.h
+++ b/src/tint/utils/text/text_style.h
@@ -31,9 +31,18 @@
 #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>
+struct ScopedTextStyle;
+}
+
 namespace tint {
 
 /// TextStyle is a styling that can be applied to span of a StyledText.
@@ -120,10 +129,37 @@
         return TextStyle{out};
     }
 
+    /// @returns a new ScopedTextStyle of @p values using with this TextStyle
+    template <typename... VALUES>
+    inline ScopedTextStyle<VALUES...> operator()(VALUES&&... values) const;
+
     /// The style bit pattern
     Bits bits = 0;
 };
 
+/// ScopedTextStyle is a span of text, styled with a TextStyle
+template <typename... VALUES>
+struct ScopedTextStyle {
+    std::tuple<VALUES...> values;
+    TextStyle style;
+};
+
+template <typename... VALUES>
+ScopedTextStyle<VALUES...> TextStyle::operator()(VALUES&&... values) const {
+    return ScopedTextStyle<VALUES...>{std::forward_as_tuple(values...), *this};
+}
+
+namespace detail {
+template <typename T>
+struct IsScopedTextStyle : std::false_type {};
+template <typename... VALUES>
+struct IsScopedTextStyle<ScopedTextStyle<VALUES...> > : std::true_type {};
+}  // namespace detail
+
+/// Resolves to true iff T is a ScopedTextStyle.
+template <typename T>
+static constexpr bool IsScopedTextStyle = detail::IsScopedTextStyle<T>::value;
+
 }  // namespace tint
 
 namespace tint::style {