diff --git a/src/tint/cmd/fuzz/ir/fuzz.h b/src/tint/cmd/fuzz/ir/fuzz.h
index 533cac4..18ffdce 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.h
+++ b/src/tint/cmd/fuzz/ir/fuzz.h
@@ -33,6 +33,7 @@
 #include <tuple>
 #include <utility>
 
+#include "src/tint/utils/bytes/buffer_reader.h"
 #include "src/tint/utils/bytes/decoder.h"
 #include "src/tint/utils/containers/slice.h"
 #include "src/tint/utils/macros/static_init.h"
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.cmake b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
index d51bd4d..350a170 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.cmake
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
@@ -55,6 +55,7 @@
   tint_lang_wgsl_program
   tint_lang_wgsl_program_fuzz
   tint_lang_wgsl_sem
+  tint_lang_wgsl_writer_raise_fuzz
   tint_lang_wgsl_fuzz
   tint_utils_bytes
   tint_utils_cli
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.gn b/src/tint/cmd/fuzz/wgsl/BUILD.gn
index 389a98e..89d7ea9 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.gn
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.gn
@@ -98,6 +98,7 @@
       "${tint_src_dir}/lang/wgsl/program",
       "${tint_src_dir}/lang/wgsl/program:fuzz",
       "${tint_src_dir}/lang/wgsl/sem",
+      "${tint_src_dir}/lang/wgsl/writer/raise:fuzz",
       "${tint_src_dir}/utils/bytes",
       "${tint_src_dir}/utils/cli",
       "${tint_src_dir}/utils/containers",
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.cc b/src/tint/cmd/fuzz/wgsl/fuzz.cc
index 6e4d9b8..23bbe0f 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.cc
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.cc
@@ -30,6 +30,8 @@
 #include <iostream>
 #include <thread>
 
+#include "src/tint/lang/wgsl/common/allowed_features.h"
+#include "src/tint/lang/wgsl/reader/options.h"
 #include "src/tint/lang/wgsl/reader/reader.h"
 #include "src/tint/utils/containers/vector.h"
 #include "src/tint/utils/macros/defer.h"
@@ -87,8 +89,14 @@
     tint::Source::File file("test.wgsl", wgsl);
 
     // Parse the WGSL program.
-    auto program = tint::wgsl::reader::Parse(&file);
+    tint::wgsl::reader::Options parse_options;
+    parse_options.allowed_features = tint::wgsl::AllowedFeatures::Everything();
+    auto program = tint::wgsl::reader::Parse(&file, parse_options);
     if (!program.IsValid()) {
+        if (options.verbose) {
+            std::cerr << "invalid WGSL program: " << std::endl
+                      << program.Diagnostics() << std::endl;
+        }
         return;
     }
 
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.h b/src/tint/cmd/fuzz/wgsl/fuzz.h
index c01455b..e4bc7a0 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.h
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.h
@@ -33,6 +33,7 @@
 #include <utility>
 
 #include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
 #include "src/tint/utils/bytes/decoder.h"
 #include "src/tint/utils/containers/slice.h"
 #include "src/tint/utils/macros/static_init.h"
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 825be15..5fbb84d 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -196,10 +196,13 @@
                 return Function::PipelineStage::kFragment;
             case pb::PipelineStage::Vertex:
                 return Function::PipelineStage::kVertex;
-            default:
-                TINT_ICE() << "unhandled PipelineStage: " << stage;
-                return Function::PipelineStage::kCompute;
+
+            case pb::PipelineStage::PipelineStage_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::PipelineStage::PipelineStage_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "unhandled PipelineStage: " << stage;
+        return Function::PipelineStage::kCompute;
     }
 
     ////////////////////////////////////////////////////////////////////////////
@@ -322,8 +325,7 @@
             case pb::Instruction::KindCase::kUnreachable:
                 inst_out = CreateInstructionUnreachable(inst_in.unreachable());
                 break;
-            default:
-                TINT_UNIMPLEMENTED() << inst_in.kind_case();
+            case pb::Instruction::KindCase::KIND_NOT_SET:
                 break;
         }
         TINT_ASSERT_OR_RETURN_VALUE(inst_out, nullptr);
@@ -550,7 +552,7 @@
                 return CreateTypeExternalTexture(type_in.external_texture());
             case pb::Type::KindCase::kSampler:
                 return CreateTypeSampler(type_in.sampler());
-            default:
+            case pb::Type::KindCase::KIND_NOT_SET:
                 break;
         }
         TINT_ICE() << type_in.kind_case();
@@ -571,10 +573,13 @@
                 return mod_out_.Types().Get<f32>();
             case pb::TypeBasic::f16:
                 return mod_out_.Types().Get<f16>();
-            default:
-                TINT_ICE() << "invalid TypeBasic: " << basic_in;
-                return nullptr;
+
+            case pb::TypeBasic::TypeBasic_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::TypeBasic::TypeBasic_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "invalid TypeBasic: " << basic_in;
+        return nullptr;
     }
 
     const type::Vector* CreateTypeVector(const pb::TypeVector& vector_in) {
@@ -712,10 +717,15 @@
             case pb::Value::KindCase::kConstant:
                 value_out = b.Constant(ConstantValue(value_in.constant()));
                 break;
-            default:
-                TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
-                return nullptr;
+            case pb::Value::KindCase::KIND_NOT_SET:
+                break;
         }
+
+        if (!value_out) {
+            TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
+            return nullptr;
+        }
+
         return value_out;
     }
 
@@ -788,10 +798,11 @@
                 return CreateConstantComposite(value_in.composite());
             case pb::ConstantValue::KindCase::kSplat:
                 return CreateConstantSplat(value_in.splat());
-            default:
-                TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
-                return nullptr;
+            case pb::ConstantValue::KindCase::KIND_NOT_SET:
+                break;
         }
+        TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
+        return nullptr;
     }
 
     const core::constant::Value* CreateConstantScalar(const pb::ConstantValueScalar& value_in) {
@@ -806,10 +817,11 @@
                 return b.ConstantValue(f32(value_in.f32()));
             case pb::ConstantValueScalar::KindCase::kF16:
                 return b.ConstantValue(f16(value_in.f16()));
-            default:
-                TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
-                return nullptr;
+            case pb::ConstantValueScalar::KindCase::KIND_NOT_SET:
+                break;
         }
+        TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
+        return nullptr;
     }
 
     const core::constant::Value* CreateConstantComposite(
@@ -874,10 +886,13 @@
                 return core::AddressSpace::kUniform;
             case pb::AddressSpace::workgroup:
                 return core::AddressSpace::kWorkgroup;
-            default:
-                TINT_ICE() << "invalid AddressSpace: " << in;
-                return core::AddressSpace::kUndefined;
+
+            case pb::AddressSpace::AddressSpace_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::AddressSpace::AddressSpace_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "invalid AddressSpace: " << in;
+        return core::AddressSpace::kUndefined;
     }
 
     core::Access AccessControl(pb::AccessControl in) {
@@ -888,10 +903,13 @@
                 return core::Access::kWrite;
             case pb::AccessControl::read_write:
                 return core::Access::kReadWrite;
-            default:
-                TINT_ICE() << "invalid Access: " << in;
-                return core::Access::kUndefined;
+
+            case pb::AccessControl::AccessControl_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::AccessControl::AccessControl_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "invalid Access: " << in;
+        return core::Access::kUndefined;
     }
 
     core::UnaryOp UnaryOp(pb::UnaryOp in) {
@@ -907,10 +925,12 @@
             case pb::UnaryOp::not_:
                 return core::UnaryOp::kNot;
 
-            default:
-                TINT_ICE() << "invalid UnaryOp: " << in;
-                return core::UnaryOp::kComplement;
+            case pb::UnaryOp::UnaryOp_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::UnaryOp::UnaryOp_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "invalid UnaryOp: " << in;
+        return core::UnaryOp::kComplement;
     }
 
     core::BinaryOp BinaryOp(pb::BinaryOp in) {
@@ -947,11 +967,17 @@
                 return core::BinaryOp::kShiftLeft;
             case pb::BinaryOp::shift_right:
                 return core::BinaryOp::kShiftRight;
+            case pb::BinaryOp::logical_and:
+                return core::BinaryOp::kLogicalAnd;
+            case pb::BinaryOp::logical_or:
+                return core::BinaryOp::kLogicalOr;
 
-            default:
-                TINT_ICE() << "invalid BinaryOp: " << in;
-                return core::BinaryOp::kAdd;
+            case pb::BinaryOp::BinaryOp_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::BinaryOp::BinaryOp_INT_MAX_SENTINEL_DO_NOT_USE_:
+                break;
         }
+        TINT_ICE() << "invalid BinaryOp: " << in;
+        return core::BinaryOp::kAdd;
     }
 
     core::type::TextureDimension TextureDimension(pb::TextureDimension in) {
@@ -968,7 +994,9 @@
                 return core::type::TextureDimension::kCube;
             case pb::TextureDimension::cube_array:
                 return core::type::TextureDimension::kCubeArray;
-            default:
+
+            case pb::TextureDimension::TextureDimension_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::TextureDimension::TextureDimension_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
 
@@ -980,6 +1008,8 @@
         switch (in) {
             case pb::TexelFormat::bgra8_unorm:
                 return core::TexelFormat::kBgra8Unorm;
+            case pb::TexelFormat::r8_unorm:
+                return core::TexelFormat::kR8Unorm;
             case pb::TexelFormat::r32_float:
                 return core::TexelFormat::kR32Float;
             case pb::TexelFormat::r32_sint:
@@ -1012,7 +1042,9 @@
                 return core::TexelFormat::kRgba8Uint;
             case pb::TexelFormat::rgba8_unorm:
                 return core::TexelFormat::kRgba8Unorm;
-            default:
+
+            case pb::TexelFormat::TexelFormat_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::TexelFormat::TexelFormat_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
 
@@ -1026,7 +1058,9 @@
                 return core::type::SamplerKind::kSampler;
             case pb::SamplerKind::comparison:
                 return core::type::SamplerKind::kComparisonSampler;
-            default:
+
+            case pb::SamplerKind::SamplerKind_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::SamplerKind::SamplerKind_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
 
@@ -1042,7 +1076,9 @@
                 return core::InterpolationType::kLinear;
             case pb::InterpolationType::perspective:
                 return core::InterpolationType::kPerspective;
-            default:
+
+            case pb::InterpolationType::InterpolationType_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::InterpolationType::InterpolationType_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
         TINT_ICE() << "invalid InterpolationType: " << in;
@@ -1057,7 +1093,9 @@
                 return core::InterpolationSampling::kCentroid;
             case pb::InterpolationSampling::sample:
                 return core::InterpolationSampling::kSample;
-            default:
+
+            case pb::InterpolationSampling::InterpolationSampling_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::InterpolationSampling::InterpolationSampling_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
         TINT_ICE() << "invalid InterpolationSampling: " << in;
@@ -1096,7 +1134,9 @@
                 return core::BuiltinValue::kVertexIndex;
             case pb::BuiltinValue::workgroup_id:
                 return core::BuiltinValue::kWorkgroupId;
-            default:
+
+            case pb::BuiltinValue::BuiltinValue_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::BuiltinValue::BuiltinValue_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
         TINT_ICE() << "invalid BuiltinValue: " << in;
@@ -1225,6 +1265,14 @@
                 return core::BuiltinFn::kPack4X8Snorm;
             case pb::BuiltinFn::pack4x8_unorm:
                 return core::BuiltinFn::kPack4X8Unorm;
+            case pb::BuiltinFn::pack4xi8:
+                return core::BuiltinFn::kPack4XI8;
+            case pb::BuiltinFn::pack4xu8:
+                return core::BuiltinFn::kPack4XU8;
+            case pb::BuiltinFn::pack4xi8_clamp:
+                return core::BuiltinFn::kPack4XI8Clamp;
+            case pb::BuiltinFn::pack4xu8_clamp:
+                return core::BuiltinFn::kPack4XU8Clamp;
             case pb::BuiltinFn::pow:
                 return core::BuiltinFn::kPow;
             case pb::BuiltinFn::quantize_to_f16:
@@ -1275,6 +1323,10 @@
                 return core::BuiltinFn::kUnpack4X8Snorm;
             case pb::BuiltinFn::unpack4x8_unorm:
                 return core::BuiltinFn::kUnpack4X8Unorm;
+            case pb::BuiltinFn::unpack4xi8:
+                return core::BuiltinFn::kUnpack4XI8;
+            case pb::BuiltinFn::unpack4xu8:
+                return core::BuiltinFn::kUnpack4XU8;
             case pb::BuiltinFn::workgroup_barrier:
                 return core::BuiltinFn::kWorkgroupBarrier;
             case pb::BuiltinFn::texture_barrier:
@@ -1335,7 +1387,9 @@
                 return core::BuiltinFn::kSubgroupBallot;
             case pb::BuiltinFn::subgroup_broadcast:
                 return core::BuiltinFn::kSubgroupBroadcast;
-            default:
+
+            case pb::BuiltinFn::BuiltinFn_INT_MIN_SENTINEL_DO_NOT_USE_:
+            case pb::BuiltinFn::BuiltinFn_INT_MAX_SENTINEL_DO_NOT_USE_:
                 break;
         }
         TINT_ICE() << "invalid BuiltinFn: " << in;
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index ef364ce..562a568 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -162,10 +162,11 @@
                 return pb::PipelineStage::Fragment;
             case Function::PipelineStage::kVertex:
                 return pb::PipelineStage::Vertex;
-            default:
-                TINT_ICE() << "unhandled PipelineStage: " << stage;
-                return pb::PipelineStage::Compute;
+            case Function::PipelineStage::kUndefined:
+                break;
         }
+        TINT_ICE() << "unhandled PipelineStage: " << stage;
+        return pb::PipelineStage::Compute;
     }
 
     ////////////////////////////////////////////////////////////////////////////
@@ -655,10 +656,14 @@
                 return pb::AddressSpace::uniform;
             case core::AddressSpace::kWorkgroup:
                 return pb::AddressSpace::workgroup;
-            default:
-                TINT_ICE() << "invalid AddressSpace: " << in;
-                return pb::AddressSpace::function;
+
+            case core::AddressSpace::kUndefined:
+            case core::AddressSpace::kIn:
+            case core::AddressSpace::kOut:
+                break;
         }
+        TINT_ICE() << "invalid AddressSpace: " << in;
+        return pb::AddressSpace::function;
     }
 
     pb::AccessControl AccessControl(core::Access in) {
@@ -669,10 +674,11 @@
                 return pb::AccessControl::write;
             case core::Access::kReadWrite:
                 return pb::AccessControl::read_write;
-            default:
-                TINT_ICE() << "invalid Access: " << in;
-                return pb::AccessControl::read;
+            case core::Access::kUndefined:
+                break;
         }
+        TINT_ICE() << "invalid Access: " << in;
+        return pb::AccessControl::read;
     }
 
     pb::UnaryOp UnaryOp(core::UnaryOp in) {
@@ -750,7 +756,7 @@
                 return pb::TextureDimension::cube;
             case core::type::TextureDimension::kCubeArray:
                 return pb::TextureDimension::cube_array;
-            default:
+            case core::type::TextureDimension::kNone:
                 break;
         }
 
@@ -768,6 +774,8 @@
                 return pb::TexelFormat::r32_sint;
             case core::TexelFormat::kR32Uint:
                 return pb::TexelFormat::r32_uint;
+            case core::TexelFormat::kR8Unorm:
+                return pb::TexelFormat::r8_unorm;
             case core::TexelFormat::kRg32Float:
                 return pb::TexelFormat::rg32_float;
             case core::TexelFormat::kRg32Sint:
@@ -794,7 +802,7 @@
                 return pb::TexelFormat::rgba8_uint;
             case core::TexelFormat::kRgba8Unorm:
                 return pb::TexelFormat::rgba8_unorm;
-            default:
+            case core::TexelFormat::kUndefined:
                 break;
         }
 
@@ -822,7 +830,7 @@
                 return pb::InterpolationType::linear;
             case core::InterpolationType::kPerspective:
                 return pb::InterpolationType::perspective;
-            default:
+            case core::InterpolationType::kUndefined:
                 break;
         }
         TINT_ICE() << "invalid InterpolationType: " << in;
@@ -837,7 +845,7 @@
                 return pb::InterpolationSampling::centroid;
             case core::InterpolationSampling::kSample:
                 return pb::InterpolationSampling::sample;
-            default:
+            case core::InterpolationSampling::kUndefined:
                 break;
         }
         TINT_ICE() << "invalid InterpolationSampling: " << in;
@@ -876,7 +884,7 @@
                 return pb::BuiltinValue::vertex_index;
             case core::BuiltinValue::kWorkgroupId:
                 return pb::BuiltinValue::workgroup_id;
-            default:
+            case core::BuiltinValue::kUndefined:
                 break;
         }
         TINT_ICE() << "invalid BuiltinValue: " << in;
@@ -1005,6 +1013,14 @@
                 return pb::BuiltinFn::pack4x8_snorm;
             case core::BuiltinFn::kPack4X8Unorm:
                 return pb::BuiltinFn::pack4x8_unorm;
+            case core::BuiltinFn::kPack4XI8:
+                return pb::BuiltinFn::pack4xi8;
+            case core::BuiltinFn::kPack4XU8:
+                return pb::BuiltinFn::pack4xu8;
+            case core::BuiltinFn::kPack4XI8Clamp:
+                return pb::BuiltinFn::pack4xi8_clamp;
+            case core::BuiltinFn::kPack4XU8Clamp:
+                return pb::BuiltinFn::pack4xu8_clamp;
             case core::BuiltinFn::kPow:
                 return pb::BuiltinFn::pow;
             case core::BuiltinFn::kQuantizeToF16:
@@ -1055,6 +1071,10 @@
                 return pb::BuiltinFn::unpack4x8_snorm;
             case core::BuiltinFn::kUnpack4X8Unorm:
                 return pb::BuiltinFn::unpack4x8_unorm;
+            case core::BuiltinFn::kUnpack4XI8:
+                return pb::BuiltinFn::unpack4xi8;
+            case core::BuiltinFn::kUnpack4XU8:
+                return pb::BuiltinFn::unpack4xu8;
             case core::BuiltinFn::kWorkgroupBarrier:
                 return pb::BuiltinFn::workgroup_barrier;
             case core::BuiltinFn::kTextureBarrier:
@@ -1115,7 +1135,7 @@
                 return pb::BuiltinFn::subgroup_ballot;
             case core::BuiltinFn::kSubgroupBroadcast:
                 return pb::BuiltinFn::subgroup_broadcast;
-            default:
+            case core::BuiltinFn::kNone:
                 break;
         }
         TINT_ICE() << "invalid BuiltinFn: " << in;
diff --git a/src/tint/lang/core/ir/binary/ir.proto b/src/tint/lang/core/ir/binary/ir.proto
index e59a9cb..19f3b0a 100644
--- a/src/tint/lang/core/ir/binary/ir.proto
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -247,31 +247,30 @@
         InstructionReturn return = 3;
         InstructionUnary unary = 4;
         InstructionBinary binary = 5;
-        InstructionBuiltin builtin = 6;
-        InstructionDiscard discard = 7;
-        InstructionLet let = 8;
-        InstructionVar var = 9;
-        InstructionBitcast bitcast = 10;
-        InstructionConstruct construct = 11;
-        InstructionConvert convert = 12;
-        InstructionAccess access = 13;
-        InstructionUserCall user_call = 14;
-        InstructionBuiltinCall builtin_call = 15;
-        InstructionLoad load = 16;
-        InstructionStore store = 17;
-        InstructionLoadVectorElement load_vector_element = 18;
-        InstructionStoreVectorElement store_vector_element = 19;
-        InstructionSwizzle swizzle = 20;
-        InstructionIf if = 21;
-        InstructionSwitch switch = 22;
-        InstructionLoop loop = 23;
-        InstructionExitIf exit_if = 24;
-        InstructionExitSwitch exit_switch = 25;
-        InstructionExitLoop exit_loop = 26;
-        InstructionNextIteration next_iteration = 27;
-        InstructionContinue continue = 28;
-        InstructionBreakIf break_if = 29;
-        InstructionUnreachable unreachable = 30;
+        InstructionDiscard discard = 6;
+        InstructionLet let = 7;
+        InstructionVar var = 8;
+        InstructionBitcast bitcast = 9;
+        InstructionConstruct construct = 10;
+        InstructionConvert convert = 11;
+        InstructionAccess access = 12;
+        InstructionUserCall user_call = 13;
+        InstructionBuiltinCall builtin_call = 14;
+        InstructionLoad load = 15;
+        InstructionStore store = 16;
+        InstructionLoadVectorElement load_vector_element = 17;
+        InstructionStoreVectorElement store_vector_element = 18;
+        InstructionSwizzle swizzle = 19;
+        InstructionIf if = 20;
+        InstructionSwitch switch = 21;
+        InstructionLoop loop = 22;
+        InstructionExitIf exit_if = 23;
+        InstructionExitSwitch exit_switch = 24;
+        InstructionExitLoop exit_loop = 25;
+        InstructionNextIteration next_iteration = 26;
+        InstructionContinue continue = 27;
+        InstructionBreakIf break_if = 28;
+        InstructionUnreachable unreachable = 29;
     }
 }
 
@@ -285,8 +284,6 @@
     BinaryOp op = 1;
 }
 
-message InstructionBuiltin {}
-
 message InstructionBitcast {}
 
 message InstructionConstruct {}
@@ -450,22 +447,23 @@
 
 enum TexelFormat {
     bgra8_unorm = 0;
-    r32_float = 1;
-    r32_sint = 2;
-    r32_uint = 3;
-    rg32_float = 4;
-    rg32_sint = 5;
-    rg32_uint = 6;
-    rgba16_float = 7;
-    rgba16_sint = 8;
-    rgba16_uint = 9;
-    rgba32_float = 10;
-    rgba32_sint = 11;
-    rgba32_uint = 12;
-    rgba8_sint = 13;
-    rgba8_snorm = 14;
-    rgba8_uint = 15;
-    rgba8_unorm = 16;
+    r8_unorm = 1;
+    r32_float = 2;
+    r32_sint = 3;
+    r32_uint = 4;
+    rg32_float = 5;
+    rg32_sint = 6;
+    rg32_uint = 7;
+    rgba16_float = 8;
+    rgba16_sint = 9;
+    rgba16_uint = 10;
+    rgba32_float = 11;
+    rgba32_sint = 12;
+    rgba32_uint = 13;
+    rgba8_sint = 14;
+    rgba8_snorm = 15;
+    rgba8_uint = 16;
+    rgba8_unorm = 17;
 }
 
 enum SamplerKind {
@@ -564,59 +562,65 @@
     pack2x16_unorm = 57;
     pack4x8_snorm = 58;
     pack4x8_unorm = 59;
-    pow = 60;
-    quantize_to_f16 = 61;
-    radians = 62;
-    reflect = 63;
-    refract = 64;
-    reverse_bits = 65;
-    round = 66;
-    saturate = 67;
-    select = 68;
-    sign = 69;
-    sin = 70;
-    sinh = 71;
-    smoothstep = 72;
-    sqrt = 73;
-    step = 74;
-    storage_barrier = 75;
-    tan = 76;
-    tanh = 77;
-    transpose = 78;
-    trunc = 79;
-    unpack2x16_float = 80;
-    unpack2x16_snorm = 81;
-    unpack2x16_unorm = 82;
-    unpack4x8_snorm = 83;
-    unpack4x8_unorm = 84;
-    workgroup_barrier = 85;
-    texture_barrier = 86;
-    texture_dimensions = 87;
-    texture_gather = 88;
-    texture_gather_compare = 89;
-    texture_num_layers = 90;
-    texture_num_levels = 91;
-    texture_num_samples = 92;
-    texture_sample = 93;
-    texture_sample_bias = 94;
-    texture_sample_compare = 95;
-    texture_sample_compare_level = 96;
-    texture_sample_grad = 97;
-    texture_sample_level = 98;
-    texture_sample_base_clamp_to_edge = 99;
-    texture_store = 100;
-    texture_load = 101;
-    atomic_load = 102;
-    atomic_store = 103;
-    atomic_add = 104;
-    atomic_sub = 105;
-    atomic_max = 106;
-    atomic_min = 107;
-    atomic_and = 108;
-    atomic_or = 109;
-    atomic_xor = 110;
-    atomic_exchange = 111;
-    atomic_compare_exchange_weak = 112;
-    subgroup_ballot = 113;
-    subgroup_broadcast = 114;
+    pack4xi8 = 60;
+    pack4xu8 = 61;
+    pack4xi8_clamp = 62;
+    pack4xu8_clamp = 63;
+    pow = 64;
+    quantize_to_f16 = 65;
+    radians = 66;
+    reflect = 67;
+    refract = 68;
+    reverse_bits = 69;
+    round = 70;
+    saturate = 71;
+    select = 72;
+    sign = 73;
+    sin = 74;
+    sinh = 75;
+    smoothstep = 76;
+    sqrt = 77;
+    step = 78;
+    storage_barrier = 79;
+    tan = 80;
+    tanh = 81;
+    transpose = 82;
+    trunc = 83;
+    unpack2x16_float = 84;
+    unpack2x16_snorm = 85;
+    unpack2x16_unorm = 86;
+    unpack4x8_snorm = 87;
+    unpack4x8_unorm = 88;
+    unpack4xi8 = 89;
+    unpack4xu8 = 90;
+    workgroup_barrier = 91;
+    texture_barrier = 92;
+    texture_dimensions = 93;
+    texture_gather = 94;
+    texture_gather_compare = 95;
+    texture_num_layers = 96;
+    texture_num_levels = 97;
+    texture_num_samples = 98;
+    texture_sample = 99;
+    texture_sample_bias = 100;
+    texture_sample_compare = 101;
+    texture_sample_compare_level = 102;
+    texture_sample_grad = 103;
+    texture_sample_level = 104;
+    texture_sample_base_clamp_to_edge = 105;
+    texture_store = 106;
+    texture_load = 107;
+    atomic_load = 108;
+    atomic_store = 109;
+    atomic_add = 110;
+    atomic_sub = 111;
+    atomic_max = 112;
+    atomic_min = 113;
+    atomic_and = 114;
+    atomic_or = 115;
+    atomic_xor = 116;
+    atomic_exchange = 117;
+    atomic_compare_exchange_weak = 118;
+    subgroup_ballot = 119;
+    subgroup_broadcast = 120;
 }
diff --git a/src/tint/lang/core/ir/block_param.h b/src/tint/lang/core/ir/block_param.h
index 009e022..0b697a2 100644
--- a/src/tint/lang/core/ir/block_param.h
+++ b/src/tint/lang/core/ir/block_param.h
@@ -31,25 +31,42 @@
 #include "src/tint/lang/core/ir/value.h"
 #include "src/tint/utils/rtti/castable.h"
 
+// Forward declarations
+namespace tint::core::ir {
+class MultiInBlock;
+}  // namespace tint::core::ir
+
 namespace tint::core::ir {
 
-/// An instruction in the IR.
+/// A block parameter in the IR.
 class BlockParam : public Castable<BlockParam, Value> {
   public:
     /// Constructor
-    /// @param type the type of the var
+    /// @param type the type of the parameter
     explicit BlockParam(const core::type::Type* type);
     ~BlockParam() override;
 
-    /// @returns the type of the var
+    /// @returns the type of the parameter
     const core::type::Type* Type() const override { return type_; }
 
+    /// Sets the block that this parameter belongs to.
+    /// @param block the block
+    void SetBlock(MultiInBlock* block) { block_ = block; }
+
+    /// @returns the block that this parameter belongs to, or nullptr
+    MultiInBlock* Block() { return block_; }
+
+    /// @returns the block that this parameter belongs to, or nullptr
+    const MultiInBlock* Block() const { return block_; }
+
     /// @copydoc Instruction::Clone()
     BlockParam* Clone(CloneContext& ctx) override;
 
   private:
-    /// the result type of the instruction
+    /// the type of the parameter
     const core::type::Type* type_ = nullptr;
+    /// the block that the parameter belongs to
+    MultiInBlock* block_ = nullptr;
 };
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index c705d0c..99e94d1 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -219,6 +219,37 @@
         cb();
     }
 
+    /// Calls @p cb with the builder inserting after @p val
+    /// @param val the insertion point for new instructions
+    /// @param cb the function to call with the builder inserting new instructions after @p val
+    template <typename FUNCTION>
+    void InsertAfter(ir::Value* val, FUNCTION&& cb) {
+        tint::Switch(
+            val,
+            [&](core::ir::InstructionResult* result) {
+                const TINT_SCOPED_ASSIGNMENT(insertion_point_,
+                                             InsertionPoints::InsertAfter{result->Instruction()});
+                cb();
+            },
+            [&](core::ir::FunctionParam* param) {
+                auto* body = param->Function()->Block();
+                if (body->IsEmpty()) {
+                    Append(body, cb);
+                } else {
+                    InsertBefore(body->Front(), cb);
+                }
+            },
+            [&](core::ir::BlockParam* param) {
+                auto* block = param->Block();
+                if (block->IsEmpty()) {
+                    Append(block, cb);
+                } else {
+                    InsertBefore(block->Front(), cb);
+                }
+            },
+            TINT_ICE_ON_NO_MATCH);
+    }
+
     /// Adds and returns the instruction @p instruction to the current insertion point. If there
     /// is no current insertion point set, then @p instruction is just returned.
     /// @param instruction the instruction to append
diff --git a/src/tint/lang/core/ir/disassembler.cc b/src/tint/lang/core/ir/disassembler.cc
index 2e95c2c..281cb61 100644
--- a/src/tint/lang/core/ir/disassembler.cc
+++ b/src/tint/lang/core/ir/disassembler.cc
@@ -188,7 +188,14 @@
     if (auto* merge = blk->As<MultiInBlock>()) {
         if (!merge->Params().IsEmpty()) {
             out_ << " (";
-            EmitValueList(merge->Params().Slice());
+            for (auto* p : merge->Params()) {
+                if (p != merge->Params().Front()) {
+                    out_ << ", ";
+                }
+                SourceMarker psm(this);
+                EmitValue(p);
+                psm.Store(p);
+            }
             out_ << ")";
         }
     }
@@ -302,7 +309,12 @@
     in_function_ = true;
 
     std::string fn_id = IdOf(func);
-    Indent() << "%" << fn_id << " =";
+    {
+        SourceMarker sm(this);
+        Indent() << "%" << fn_id;
+        sm.Store(func);
+    }
+    out_ << " =";
 
     if (func->Stage() != Function::PipelineStage::kUndefined) {
         out_ << " @" << func->Stage();
@@ -318,7 +330,9 @@
         if (p != func->Params().Front()) {
             out_ << ", ";
         }
+        SourceMarker sm(this);
         out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+        sm.Store(p);
 
         EmitParamAttributes(p);
     }
@@ -790,15 +804,6 @@
     );
 }
 
-void Disassembler::EmitValueList(tint::Slice<const Value* const> values) {
-    for (size_t i = 0, n = values.Length(); i < n; i++) {
-        if (i > 0) {
-            out_ << ", ";
-        }
-        EmitValue(values[i]);
-    }
-}
-
 void Disassembler::EmitBinary(const Binary* b) {
     SourceMarker sm(this);
     EmitValueWithType(b);
diff --git a/src/tint/lang/core/ir/disassembler.h b/src/tint/lang/core/ir/disassembler.h
index d22b8a1..6b6f7c6 100644
--- a/src/tint/lang/core/ir/disassembler.h
+++ b/src/tint/lang/core/ir/disassembler.h
@@ -32,6 +32,7 @@
 
 #include "src/tint/lang/core/ir/binary.h"
 #include "src/tint/lang/core/ir/block.h"
+#include "src/tint/lang/core/ir/block_param.h"
 #include "src/tint/lang/core/ir/call.h"
 #include "src/tint/lang/core/ir/if.h"
 #include "src/tint/lang/core/ir/loop.h"
@@ -100,10 +101,26 @@
     /// @returns the source for the result
     Source ResultSource(IndexedValue result) { return result_to_src_.GetOr(result, Source{}); }
 
-    /// @param blk teh block to retrieve
+    /// @param blk the block to retrieve
     /// @returns the source for the block
     Source BlockSource(const Block* blk) { return block_to_src_.GetOr(blk, Source{}); }
 
+    /// @param param the block parameter to retrieve
+    /// @returns the source for the parameter
+    Source BlockParamSource(const BlockParam* param) {
+        return block_param_to_src_.GetOr(param, Source{});
+    }
+
+    /// @param func the function to retrieve
+    /// @returns the source for the function
+    Source FunctionSource(const Function* func) { return function_to_src_.GetOr(func, Source{}); }
+
+    /// @param param the function parameter to retrieve
+    /// @returns the source for the parameter
+    Source FunctionParamSource(const FunctionParam* param) {
+        return function_param_to_src_.GetOr(param, Source{});
+    }
+
     /// Stores the given @p src location for @p inst instruction
     /// @param inst the instruction to store
     /// @param src the source location
@@ -114,6 +131,23 @@
     /// @param src the source location
     void SetSource(const Block* blk, Source src) { block_to_src_.Add(blk, src); }
 
+    /// Stores the given @p src location for @p param block parameter
+    /// @param param the block parameter to store
+    /// @param src the source location
+    void SetSource(const BlockParam* param, Source src) { block_param_to_src_.Add(param, src); }
+
+    /// Stores the given @p src location for @p func function
+    /// @param func the function to store
+    /// @param src the source location
+    void SetSource(const Function* func, Source src) { function_to_src_.Add(func, src); }
+
+    /// Stores the given @p src location for @p param function parameter
+    /// @param param the function parameter to store
+    /// @param src the source location
+    void SetSource(const FunctionParam* param, Source src) {
+        function_param_to_src_.Add(param, src);
+    }
+
     /// Stores the given @p src location for @p op operand
     /// @param op the operand to store
     /// @param src the source location
@@ -137,6 +171,12 @@
 
         void Store(const Block* blk) { dis_->SetSource(blk, MakeSource()); }
 
+        void Store(const BlockParam* param) { dis_->SetSource(param, MakeSource()); }
+
+        void Store(const Function* func) { dis_->SetSource(func, MakeSource()); }
+
+        void Store(const FunctionParam* param) { dis_->SetSource(param, MakeSource()); }
+
         void Store(IndexedValue operand) { dis_->SetSource(operand, MakeSource()); }
 
         void StoreResult(IndexedValue result) { dis_->SetResultSource(result, MakeSource()); }
@@ -168,7 +208,6 @@
     void EmitValueWithType(const Instruction* val);
     void EmitValueWithType(const Value* val);
     void EmitValue(const Value* val);
-    void EmitValueList(tint::Slice<const ir::Value* const> values);
     void EmitBinary(const Binary* b);
     void EmitUnary(const Unary* b);
     void EmitTerminator(const Terminator* b);
@@ -193,9 +232,12 @@
     uint32_t current_output_start_pos_ = 0;
 
     Hashmap<const Block*, Source, 8> block_to_src_;
+    Hashmap<const BlockParam*, Source, 8> block_param_to_src_;
     Hashmap<const Instruction*, Source, 8> instruction_to_src_;
     Hashmap<IndexedValue, Source, 8> operand_to_src_;
     Hashmap<IndexedValue, Source, 8> result_to_src_;
+    Hashmap<const Function*, Source, 8> function_to_src_;
+    Hashmap<const FunctionParam*, Source, 8> function_param_to_src_;
     Hashmap<const If*, std::string, 8> if_names_;
     Hashmap<const Loop*, std::string, 8> loop_names_;
     Hashmap<const Switch*, std::string, 8> switch_names_;
diff --git a/src/tint/lang/core/ir/function.cc b/src/tint/lang/core/ir/function.cc
index 51f7928..5b669c2 100644
--- a/src/tint/lang/core/ir/function.cc
+++ b/src/tint/lang/core/ir/function.cc
@@ -53,7 +53,7 @@
     auto* new_func =
         ctx.ir.allocators.values.Create<Function>(return_.type, pipeline_stage_, workgroup_size_);
     new_func->block_ = ctx.ir.blocks.Create<ir::Block>();
-    new_func->params_ = ctx.Clone<1>(params_.Slice());
+    new_func->SetParams(ctx.Clone<1>(params_.Slice()));
     new_func->return_.builtin = return_.builtin;
     new_func->return_.location = return_.location;
     new_func->return_.invariant = return_.invariant;
@@ -66,13 +66,25 @@
 }
 
 void Function::SetParams(VectorRef<FunctionParam*> params) {
+    for (auto* param : params_) {
+        param->SetFunction(nullptr);
+    }
     params_ = std::move(params);
-    TINT_ASSERT(!params_.Any(IsNull));
+    TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+    for (auto* param : params_) {
+        param->SetFunction(this);
+    }
 }
 
 void Function::SetParams(std::initializer_list<FunctionParam*> params) {
+    for (auto* param : params_) {
+        param->SetFunction(nullptr);
+    }
     params_ = params;
-    TINT_ASSERT(!params_.Any(IsNull));
+    TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+    for (auto* param : params_) {
+        param->SetFunction(this);
+    }
 }
 
 void Function::Destroy() {
diff --git a/src/tint/lang/core/ir/function_param.h b/src/tint/lang/core/ir/function_param.h
index 6cefb10..f323914 100644
--- a/src/tint/lang/core/ir/function_param.h
+++ b/src/tint/lang/core/ir/function_param.h
@@ -38,6 +38,11 @@
 #include "src/tint/utils/ice/ice.h"
 #include "src/tint/utils/rtti/castable.h"
 
+// Forward declarations
+namespace tint::core::ir {
+class Function;
+}  // namespace tint::core::ir
+
 namespace tint::core::ir {
 
 /// A function parameter in the IR.
@@ -48,6 +53,16 @@
     explicit FunctionParam(const core::type::Type* type);
     ~FunctionParam() override;
 
+    /// Sets the function that this parameter belongs to.
+    /// @param func the function
+    void SetFunction(ir::Function* func) { func_ = func; }
+
+    /// @returns the function that this parameter belongs to, or nullptr
+    ir::Function* Function() { return func_; }
+
+    /// @returns the function that this parameter belongs to, or nullptr
+    const ir::Function* Function() const { return func_; }
+
     /// @returns the type of the var
     const core::type::Type* Type() const override { return type_; }
 
@@ -99,6 +114,7 @@
     std::optional<struct BindingPoint> BindingPoint() const { return binding_point_; }
 
   private:
+    ir::Function* func_ = nullptr;
     const core::type::Type* type_ = nullptr;
     std::optional<core::BuiltinValue> builtin_;
     std::optional<struct Location> location_;
diff --git a/src/tint/lang/core/ir/function_test.cc b/src/tint/lang/core/ir/function_test.cc
index 8873648..27f332e 100644
--- a/src/tint/lang/core/ir/function_test.cc
+++ b/src/tint/lang/core/ir/function_test.cc
@@ -126,6 +126,12 @@
 
     // Cloned functions are not automatically added to the module.
     EXPECT_EQ(mod.functions.Length(), 1u);
+
+    // Check parameter ownership is correct.
+    EXPECT_EQ(param1->Function(), f);
+    EXPECT_EQ(param2->Function(), f);
+    EXPECT_EQ(new_param1->Function(), new_f);
+    EXPECT_EQ(new_param2->Function(), new_f);
 }
 
 TEST_F(IR_FunctionTest, CloneWithExits) {
@@ -138,5 +144,23 @@
     EXPECT_EQ(new_f, new_f->Block()->Front()->As<Return>()->Func());
 }
 
+TEST_F(IR_FunctionTest, Parameters) {
+    auto* f = b.Function("my_func", mod.Types().void_());
+
+    auto* param1 = b.FunctionParam("a", mod.Types().i32());
+    auto* param2 = b.FunctionParam("b", mod.Types().f32());
+    auto* param3 = b.FunctionParam("b", mod.Types().f32());
+
+    f->SetParams({param1, param2});
+    EXPECT_EQ(param1->Function(), f);
+    EXPECT_EQ(param2->Function(), f);
+    EXPECT_EQ(param3->Function(), nullptr);
+
+    f->SetParams({param1, param3});
+    EXPECT_EQ(param1->Function(), f);
+    EXPECT_EQ(param2->Function(), nullptr);
+    EXPECT_EQ(param3->Function(), f);
+}
+
 }  // namespace
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/multi_in_block.cc b/src/tint/lang/core/ir/multi_in_block.cc
index c57c622..411d06c 100644
--- a/src/tint/lang/core/ir/multi_in_block.cc
+++ b/src/tint/lang/core/ir/multi_in_block.cc
@@ -55,11 +55,25 @@
 }
 
 void MultiInBlock::SetParams(VectorRef<BlockParam*> params) {
+    for (auto* param : params_) {
+        param->SetBlock(nullptr);
+    }
     params_ = std::move(params);
+    TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+    for (auto* param : params_) {
+        param->SetBlock(this);
+    }
 }
 
 void MultiInBlock::SetParams(std::initializer_list<BlockParam*> params) {
+    for (auto* param : params_) {
+        param->SetBlock(nullptr);
+    }
     params_ = std::move(params);
+    TINT_ASSERT_OR_RETURN(!params_.Any(IsNull));
+    for (auto* param : params_) {
+        param->SetBlock(this);
+    }
 }
 
 void MultiInBlock::AddInboundSiblingBranch(ir::Terminator* node) {
diff --git a/src/tint/lang/core/ir/multi_in_block_test.cc b/src/tint/lang/core/ir/multi_in_block_test.cc
index 790a90e..a377c44 100644
--- a/src/tint/lang/core/ir/multi_in_block_test.cc
+++ b/src/tint/lang/core/ir/multi_in_block_test.cc
@@ -54,7 +54,9 @@
     auto* blk = b.MultiInBlock();
     auto* add = b.Add(mod.Types().i32(), 1_i, 2_i);
     blk->Append(add);
-    blk->SetParams({b.BlockParam(mod.Types().i32()), b.BlockParam(mod.Types().f32())});
+    auto* param1 = b.BlockParam(mod.Types().i32());
+    auto* param2 = b.BlockParam(mod.Types().f32());
+    blk->SetParams({param1, param2});
     blk->SetParent(loop);
 
     auto* terminate = b.TerminateInvocation();
@@ -75,6 +77,12 @@
     EXPECT_NE(add, new_blk->Front());
     EXPECT_TRUE(new_blk->Front()->Is<Binary>());
     EXPECT_EQ(BinaryOp::kAdd, new_blk->Front()->As<Binary>()->Op());
+
+    // Check parameter ownership is correct.
+    EXPECT_EQ(param1->Block(), blk);
+    EXPECT_EQ(param2->Block(), blk);
+    EXPECT_EQ(new_blk->Params()[0]->Block(), new_blk);
+    EXPECT_EQ(new_blk->Params()[1]->Block(), new_blk);
 }
 
 TEST_F(IR_MultiInBlockTest, CloneEmpty) {
@@ -86,5 +94,23 @@
     EXPECT_EQ(0u, new_blk->Params().Length());
 }
 
+TEST_F(IR_MultiInBlockTest, Parameters) {
+    auto* blk = b.MultiInBlock();
+
+    auto* param1 = b.BlockParam("a", mod.Types().i32());
+    auto* param2 = b.BlockParam("b", mod.Types().f32());
+    auto* param3 = b.BlockParam("b", mod.Types().f32());
+
+    blk->SetParams({param1, param2});
+    EXPECT_EQ(param1->Block(), blk);
+    EXPECT_EQ(param2->Block(), blk);
+    EXPECT_EQ(param3->Block(), nullptr);
+
+    blk->SetParams({param1, param3});
+    EXPECT_EQ(param1->Block(), blk);
+    EXPECT_EQ(param2->Block(), nullptr);
+    EXPECT_EQ(param3->Block(), blk);
+}
+
 }  // namespace
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 3f1a681..24d594d 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -158,11 +158,26 @@
     /// @returns the diagnostic
     diag::Diagnostic& AddResultError(const Instruction* inst, size_t idx);
 
-    /// Adds an error the @p block and highlights the block header in the disassembly
+    /// Adds an error for the @p block and highlights the block header in the disassembly
     /// @param blk the block
     /// @returns the diagnostic
     diag::Diagnostic& AddError(const Block* blk);
 
+    /// Adds an error for the @p param and highlights the parameter in the disassembly
+    /// @param param the parameter
+    /// @returns the diagnostic
+    diag::Diagnostic& AddError(const BlockParam* param);
+
+    /// Adds an error for the @p func and highlights the function in the disassembly
+    /// @param func the function
+    /// @returns the diagnostic
+    diag::Diagnostic& AddError(const Function* func);
+
+    /// Adds an error for the @p param and highlights the parameter in the disassembly
+    /// @param param the parameter
+    /// @returns the diagnostic
+    diag::Diagnostic& AddError(const FunctionParam* param);
+
     /// Adds an error the @p block and highlights the block header in the disassembly
     /// @param src the source lines to highlight
     /// @returns the diagnostic
@@ -172,6 +187,10 @@
     /// @param inst the instruction
     diag::Diagnostic& AddNote(const Instruction* inst);
 
+    /// Adds a note to @p func and highlights the function in the disassembly
+    /// @param func the function
+    diag::Diagnostic& AddNote(const Function* func);
+
     /// Adds a note to @p inst for operand @p idx and highlights the operand in the
     /// disassembly
     /// @param inst the instruction
@@ -346,8 +365,8 @@
 
     for (auto& func : mod_.functions) {
         if (!all_functions_.Add(func.Get())) {
-            AddError(Source{}) << "function " << style::Function(Name(func.Get()))
-                               << " added to module multiple times";
+            AddError(func) << "function " << style::Function(Name(func.Get()))
+                           << " added to module multiple times";
         }
     }
 
@@ -413,12 +432,36 @@
     return AddError(src);
 }
 
+diag::Diagnostic& Validator::AddError(const BlockParam* param) {
+    DisassembleIfNeeded();
+    auto src = dis_.BlockParamSource(param);
+    return AddError(src);
+}
+
+diag::Diagnostic& Validator::AddError(const Function* func) {
+    DisassembleIfNeeded();
+    auto src = dis_.FunctionSource(func);
+    return AddError(src);
+}
+
+diag::Diagnostic& Validator::AddError(const FunctionParam* param) {
+    DisassembleIfNeeded();
+    auto src = dis_.FunctionParamSource(param);
+    return AddError(src);
+}
+
 diag::Diagnostic& Validator::AddNote(const Instruction* inst) {
     DisassembleIfNeeded();
     auto src = dis_.InstructionSource(inst);
     return AddNote(src);
 }
 
+diag::Diagnostic& Validator::AddNote(const Function* func) {
+    DisassembleIfNeeded();
+    auto src = dis_.FunctionSource(func);
+    return AddNote(src);
+}
+
 diag::Diagnostic& Validator::AddNote(const Instruction* inst, size_t idx) {
     DisassembleIfNeeded();
     auto src = dis_.OperandSource(Disassembler::IndexedValue{inst, static_cast<uint32_t>(idx)});
@@ -488,22 +531,50 @@
 void Validator::CheckFunction(const Function* func) {
     CheckBlock(func->Block());
 
-    // References not allowed on function signatures even with Capability::kAllowRefTypes
     for (auto* param : func->Params()) {
+        if (!param->Alive()) {
+            AddError(param) << "destroyed parameter found in function parameter list";
+            return;
+        }
+        if (!param->Function()) {
+            AddError(param) << "function parameter has nullptr parent function";
+            return;
+        } else if (param->Function() != func) {
+            AddError(param) << "function parameter has incorrect parent function";
+            AddNote(param->Function()) << "parent function declared here";
+            return;
+        }
+
+        // References not allowed on function signatures even with Capability::kAllowRefTypes
         if (HoldsType<type::Reference>(param->Type())) {
-            // TODO(dsinclair): Parameters need a source mapping.
-            AddError(Source{}) << "references are not permitted as parameter types";
+            AddError(param) << "references are not permitted as parameter types";
         }
     }
     if (HoldsType<type::Reference>(func->ReturnType())) {
-        // TODO(dsinclair): Function need a source mapping.
-        AddError(Source{}) << "references are not permitted as return types";
+        AddError(func) << "references are not permitted as return types";
     }
 }
 
 void Validator::CheckBlock(const Block* blk) {
     TINT_SCOPED_ASSIGNMENT(current_block_, blk);
 
+    if (auto* mb = blk->As<MultiInBlock>()) {
+        for (auto* param : mb->Params()) {
+            if (!param->Alive()) {
+                AddError(param) << "destroyed parameter found in block parameter list";
+                return;
+            }
+            if (!param->Block()) {
+                AddError(param) << "block parameter has nullptr parent block";
+                return;
+            } else if (param->Block() != mb) {
+                AddError(param) << "block parameter has incorrect parent block";
+                AddNote(param->Block()) << "parent block declared here";
+                return;
+            }
+        }
+    }
+
     if (!blk->Terminator()) {
         AddError(blk) << "block: does not end in a terminator instruction";
     }
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 3e1a708..8742f8a 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -139,7 +139,10 @@
     auto res = ir::Validate(mod);
     ASSERT_NE(res, Success);
     EXPECT_EQ(res.Failure().reason.Str(),
-              R"(error: function 'my_func' added to module multiple times
+              R"(:1:1 error: function 'my_func' added to module multiple times
+%my_func = func(%2:i32, %3:f32):void -> %b1 {
+^^^^^^^^
+
 note: # Disassembly
 %my_func = func(%2:i32, %3:f32):void -> %b1 {
   %b1 = block {
@@ -154,6 +157,88 @@
 )");
 }
 
+TEST_F(IR_ValidatorTest, Function_DeadParameter) {
+    auto* f = b.Function("my_func", ty.void_());
+    auto* p = b.FunctionParam("my_param", ty.f32());
+    f->SetParams({p});
+    f->Block()->Append(b.Return(f));
+
+    p->Destroy();
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:1:17 error: destroyed parameter found in function parameter list
+%my_func = func(%my_param:f32):void -> %b1 {
+                ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:f32):void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_ParameterWithNullFunction) {
+    auto* f = b.Function("my_func", ty.void_());
+    auto* p = b.FunctionParam("my_param", ty.f32());
+    f->SetParams({p});
+    f->Block()->Append(b.Return(f));
+
+    p->SetFunction(nullptr);
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:1:17 error: function parameter has nullptr parent function
+%my_func = func(%my_param:f32):void -> %b1 {
+                ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func(%my_param:f32):void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_ParameterUsedInMultipleFunctions) {
+    auto* p = b.FunctionParam("my_param", ty.f32());
+    auto* f1 = b.Function("my_func1", ty.void_());
+    auto* f2 = b.Function("my_func2", ty.void_());
+    f1->SetParams({p});
+    f2->SetParams({p});
+    f1->Block()->Append(b.Return(f1));
+    f2->Block()->Append(b.Return(f2));
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:1:18 error: function parameter has incorrect parent function
+%my_func1 = func(%my_param:f32):void -> %b1 {
+                 ^^^^^^^^^^^^^
+
+:6:1 note: parent function declared here
+%my_func2 = func(%my_param:f32):void -> %b2 {
+^^^^^^^^^
+
+note: # Disassembly
+%my_func1 = func(%my_param:f32):void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+%my_func2 = func(%my_param:f32):void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)");
+}
+
 TEST_F(IR_ValidatorTest, CallToFunctionOutsideModule) {
     auto* f = b.Function("f", ty.void_());
     auto* g = b.Function("g", ty.void_());
@@ -393,6 +478,115 @@
 )");
 }
 
+TEST_F(IR_ValidatorTest, Block_DeadParameter) {
+    auto* f = b.Function("my_func", ty.void_());
+
+    auto* p = b.BlockParam("my_param", ty.f32());
+    b.Append(f->Block(), [&] {
+        auto* l = b.Loop();
+        l->Body()->SetParams({p});
+        b.Append(l->Body(), [&] { b.ExitLoop(l); });
+        b.Return(f);
+    });
+
+    p->Destroy();
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:4:20 error: destroyed parameter found in block parameter list
+      %b2 = block (%my_param:f32) {  # body
+                   ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%my_param:f32) {  # body
+        exit_loop  # loop_1
+      }
+    }
+    ret
+  }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Block_ParameterWithNullBlock) {
+    auto* f = b.Function("my_func", ty.void_());
+
+    auto* p = b.BlockParam("my_param", ty.f32());
+    b.Append(f->Block(), [&] {
+        auto* l = b.Loop();
+        l->Body()->SetParams({p});
+        b.Append(l->Body(), [&] { b.ExitLoop(l); });
+        b.Return(f);
+    });
+
+    p->SetBlock(nullptr);
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:4:20 error: block parameter has nullptr parent block
+      %b2 = block (%my_param:f32) {  # body
+                   ^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%my_param:f32) {  # body
+        exit_loop  # loop_1
+      }
+    }
+    ret
+  }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Block_ParameterUsedInMultipleBlocks) {
+    auto* f = b.Function("my_func", ty.void_());
+
+    auto* p = b.BlockParam("my_param", ty.f32());
+    b.Append(f->Block(), [&] {
+        auto* l = b.Loop();
+        l->Body()->SetParams({p});
+        b.Append(l->Body(), [&] { b.Continue(l, p); });
+        l->Continuing()->SetParams({p});
+        b.Append(l->Continuing(), [&] { b.NextIteration(l, p); });
+        b.Return(f);
+    });
+
+    auto res = ir::Validate(mod);
+    ASSERT_NE(res, Success);
+    EXPECT_EQ(res.Failure().reason.Str(),
+              R"(:4:20 error: block parameter has incorrect parent block
+      %b2 = block (%my_param:f32) {  # body
+                   ^^^^^^^^^^^^^
+
+:7:7 note: parent block declared here
+      %b3 = block (%my_param:f32) {  # continuing
+      ^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = func():void -> %b1 {
+  %b1 = block {
+    loop [b: %b2, c: %b3] {  # loop_1
+      %b2 = block (%my_param:f32) {  # body
+        continue %b3 %my_param:f32
+      }
+      %b3 = block (%my_param:f32) {  # continuing
+        next_iteration %b2 %my_param:f32
+      }
+    }
+    ret
+  }
+}
+)");
+}
+
 TEST_F(IR_ValidatorTest, Access_NegativeIndex) {
     auto* f = b.Function("my_func", ty.void_());
     auto* obj = b.FunctionParam(ty.vec3<f32>());
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
index 04e7b14..fd59677 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local.cc
@@ -39,7 +39,10 @@
 #include "src/tint/lang/wgsl/sem/statement.h"
 #include "src/tint/lang/wgsl/sem/struct.h"
 #include "src/tint/utils/containers/transform.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+#include "src/tint/utils/result/result.h"
 #include "src/tint/utils/rtti/switch.h"
+#include "src/tint/utils/text/text_style.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal);
 TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::PixelLocal::RasterizerOrderedView);
@@ -120,7 +123,12 @@
                 // Obtain struct of the pixel local.
                 auto* pixel_local_str =
                     pixel_local_variable->Type()->UnwrapRef()->As<sem::Struct>();
-                TransformEntryPoint(entry_point, pixel_local_variable, pixel_local_str);
+                if (auto res =
+                        TransformEntryPoint(entry_point, pixel_local_variable, pixel_local_str);
+                    res != Success) {
+                    b.Diagnostics().Add(res.Failure().reason);
+                    made_changes = true;
+                }
 
                 break;  // Only a single `var<pixel_local>` can be used by an entry point.
             }
@@ -139,9 +147,9 @@
     /// @param entry_point the entry point
     /// @param pixel_local_var the `var<pixel_local>`
     /// @param pixel_local_str the struct type of the var
-    void TransformEntryPoint(const sem::Function* entry_point,
-                             const sem::GlobalVariable* pixel_local_var,
-                             const sem::Struct* pixel_local_str) {
+    Result<SuccessType> TransformEntryPoint(const sem::Function* entry_point,
+                                            const sem::GlobalVariable* pixel_local_var,
+                                            const sem::Struct* pixel_local_str) {
         // Wrap the old entry point "fn" into a new entry point where functions to load and store
         // ROV data are called.
         auto* original_entry_point_fn = entry_point->Declaration();
@@ -175,9 +183,12 @@
         // load data from and store data into the ROVs.
         auto load_rov_function_name = b.Symbols().New("load_from_pixel_local_storage");
         auto store_rov_function_name = b.Symbols().New("store_into_pixel_local_storage");
-        DeclareROVsAndLoadStoreFunctions(load_rov_function_name, store_rov_function_name,
-                                         pixel_local_var->Declaration()->name->symbol.Name(),
-                                         pixel_local_str);
+        if (auto res = DeclareROVsAndLoadStoreFunctions(
+                load_rov_function_name, store_rov_function_name,
+                pixel_local_var->Declaration()->name->symbol.Name(), pixel_local_str);
+            res != Success) {
+            return res.Failure();
+        }
 
         // Declare new entry point
         Vector<const ast::Statement*, 5> new_entry_point_function_body;
@@ -276,6 +287,7 @@
         // Declare the new entry point that calls the inner function
         b.Func(entry_point_name, std::move(new_entry_point_params), new_entry_point_return_type,
                new_entry_point_function_body, Vector{b.Stage(ast::PipelineStage::kFragment)});
+        return Success;
     }
 
     /// Add the declarations of all the ROVs as a special type of read-write storage texture that
@@ -285,10 +297,11 @@
     /// @param store_rov_function_name the name of the function that stores the data into the ROVs
     /// @param pixel_local_variable_name the name of the pixel local variable
     /// @param pixel_local_str the struct type of the pixel local variable
-    void DeclareROVsAndLoadStoreFunctions(const Symbol& load_rov_function_name,
-                                          const Symbol& store_rov_function_name,
-                                          const std::string& pixel_local_variable_name,
-                                          const sem::Struct* pixel_local_str) {
+    Result<SuccessType> DeclareROVsAndLoadStoreFunctions(
+        const Symbol& load_rov_function_name,
+        const Symbol& store_rov_function_name,
+        const std::string& pixel_local_variable_name,
+        const sem::Struct* pixel_local_str) {
         std::string_view load_store_input_name = "my_input";
         Vector load_parameters{b.Param(load_store_input_name, b.ty.vec4<f32>())};
         Vector store_parameters{b.Param(load_store_input_name, b.ty.vec4<f32>())};
@@ -314,8 +327,11 @@
                 [&](const core::type::F32*) { return core::TexelFormat::kR32Float; },
                 TINT_ICE_ON_NO_MATCH);
             auto rov_format = ROVTexelFormat(member->Index());
-            auto rov_type = b.ty.storage_texture(core::type::TextureDimension::k2d, rov_format,
-                                                 core::Access::kReadWrite);
+            if (TINT_UNLIKELY(rov_format != Success)) {
+                return rov_format.Failure();
+            }
+            auto rov_type = b.ty.storage_texture(core::type::TextureDimension::k2d,
+                                                 rov_format.Get(), core::Access::kReadWrite);
             auto rov_symbol_name = b.Symbols().New("pixel_local_" + member->Name().Name());
             b.GlobalVar(rov_symbol_name, rov_type,
                         tint::Vector{b.Binding(AInt(ROVRegisterIndex(member->Index()))),
@@ -352,7 +368,7 @@
             // textureStore(
             //     pixel_local_member, rov_texcoord, vec4u(bitcast(PLS_Private_Variable.member)));
             std::string rov_pixel_type;
-            switch (rov_format) {
+            switch (rov_format.Get()) {
                 case core::TexelFormat::kR32Uint:
                     rov_pixel_type = "vec4u";
                     break;
@@ -363,8 +379,7 @@
                     rov_pixel_type = "vec4f";
                     break;
                 default:
-                    TINT_ICE() << "Invalid ROV format (now only R32Uint, R32Sint and R32Float are "
-                                  "supported)";
+                    TINT_UNREACHABLE();
                     break;
             }
             auto pixel_local_var_member_access_in_store_call =
@@ -374,7 +389,7 @@
                 to_vec4_call = b.Call(rov_pixel_type, pixel_local_var_member_access_in_store_call);
             } else {
                 ast::Type rov_pixel_ast_type;
-                switch (rov_format) {
+                switch (rov_format.Get()) {
                     case core::TexelFormat::kR32Uint:
                         rov_pixel_ast_type = b.ty.u32();
                         break;
@@ -399,6 +414,7 @@
 
         b.Func(load_rov_function_name, std::move(load_parameters), b.ty.void_(), load_body);
         b.Func(store_rov_function_name, std::move(store_parameters), b.ty.void_(), store_body);
+        return Success;
     }
 
     /// Find and get `@builtin(position)` which is needed for loading and storing data with ROVs
@@ -462,12 +478,14 @@
 
     /// @returns the texel format for the pixel local field with the given index
     /// @param field_index the pixel local field index
-    core::TexelFormat ROVTexelFormat(uint32_t field_index) {
+    Result<core::TexelFormat> ROVTexelFormat(uint32_t field_index) {
         auto format = cfg.pls_member_to_rov_format.Get(field_index);
         if (TINT_UNLIKELY(!format)) {
-            b.Diagnostics().AddError(diag::System::Transform, Source{})
-                << "PixelLocal::Config::attachments missing entry for field " << field_index;
-            return core::TexelFormat::kUndefined;
+            diag::Diagnostic err;
+            err.severity = diag::Severity::Error;
+            err.message << "PixelLocal::Config::attachments missing entry for field "
+                        << field_index;
+            return Failure{std::move(err)};
         }
         return *format;
     }
diff --git a/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc b/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
index 9f644a7..f34138a 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/pixel_local_test.cc
@@ -61,6 +61,31 @@
     EXPECT_FALSE(ShouldRun<PixelLocal>(src, Bindings({})));
 }
 
+TEST_F(HLSLPixelLocalTest, MissingBindings) {
+    auto* src = R"(
+enable chromium_experimental_pixel_local;
+
+struct PixelLocal {
+  a : u32,
+}
+
+var<pixel_local> P : PixelLocal;
+
+@fragment
+fn F() -> @location(0) vec4f {
+  P.a += 42;
+  return vec4f(1, 0, 0, 1);
+}
+)";
+
+    auto* expect = R"(error: PixelLocal::Config::attachments missing entry for field 0)";
+    ast::transform::DataMap data;
+    data.Add<PixelLocal::Config>();
+    auto got = Run<PixelLocal>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(HLSLPixelLocalTest, UseInEntryPoint_NoPosition) {
     auto* src = R"(
 enable chromium_experimental_pixel_local;
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
index 901b4bb..98b3067 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
@@ -134,7 +134,7 @@
     }
 
     // Replace each access instruction that we recorded.
-    Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_local;
+    Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_var;
     Hashmap<PartialAccess, core::ir::Value*, 4> source_object_to_value;
     for (const auto& to_replace : worklist) {
         auto* access = to_replace.access;
@@ -145,24 +145,58 @@
         if (to_replace.first_dynamic_index > 0) {
             PartialAccess partial_access = {
                 access->Object(), access->Indices().Truncate(to_replace.first_dynamic_index)};
-            source_object = source_object_to_value.GetOrAdd(partial_access, [&] {
-                auto* intermediate_source = builder.Access(to_replace.dynamic_index_source_type,
-                                                           source_object, partial_access.indices);
-                intermediate_source->InsertBefore(access);
-                return intermediate_source->Result(0);
-            });
+            source_object =
+                source_object_to_value.GetOrAdd(partial_access, [&]() -> core::ir::Value* {
+                    // If the source is a constant, then the partial access will also produce a
+                    // constant. Extract the constant::Value and use that as the new source object.
+                    if (source_object->Is<core::ir::Constant>()) {
+                        for (const auto& i : partial_access.indices) {
+                            auto idx = i->As<core::ir::Constant>()->Value()->ValueAs<uint32_t>();
+                            source_object = builder.Constant(
+                                source_object->As<core::ir::Constant>()->Value()->Index(idx));
+                        }
+                        return source_object;
+                    }
+
+                    // Extract a non-constant intermediate source using an access instruction that
+                    // we insert immediately after the definition of the root source object.
+                    auto* intermediate_source =
+                        builder.Access(to_replace.dynamic_index_source_type, source_object,
+                                       partial_access.indices);
+                    builder.InsertAfter(source_object,
+                                        [&] { builder.Append(intermediate_source); });
+                    return intermediate_source->Result(0);
+                });
         }
 
-        // Declare a local variable and copy the source object to it.
-        auto* local = object_to_local.GetOrAdd(source_object, [&] {
-            auto* decl = builder.Var(ir.Types().ptr(
-                core::AddressSpace::kFunction, source_object->Type(), core::Access::kReadWrite));
+        // Declare a variable and copy the source object to it.
+        auto* var = object_to_var.GetOrAdd(source_object, [&] {
+            // If the source object is a constant we use a module-scope variable, as it could be
+            // indexed by multiple functions. Otherwise, we declare a function-scope variable
+            // immediately after the definition of the source object.
+            core::ir::Var* decl = nullptr;
+            if (source_object->Is<core::ir::Constant>()) {
+                decl = builder.Var(ir.Types().ptr(core::AddressSpace::kPrivate,
+                                                  source_object->Type(), core::Access::kReadWrite));
+                ir.root_block->Append(decl);
+            } else {
+                builder.InsertAfter(source_object, [&] {
+                    decl = builder.Var(ir.Types().ptr(core::AddressSpace::kFunction,
+                                                      source_object->Type(),
+                                                      core::Access::kReadWrite));
+
+                    // If we ever support value declarations at module-scope, we will need to modify
+                    // the partial access logic above since `access` instructions cannot be used in
+                    // the root block.
+                    TINT_ASSERT(decl->Block() != ir.root_block);
+                });
+            }
+
             decl->SetInitializer(source_object);
-            decl->InsertBefore(access);
             return decl->Result(0);
         });
 
-        // Create a new access instruction using the local variable as the source.
+        // Create a new access instruction using the new variable as the source.
         Vector<core::ir::Value*, 4> indices{
             access->Indices().Offset(to_replace.first_dynamic_index)};
         const core::type::Type* access_type = access->Result(0)->Type();
@@ -178,9 +212,9 @@
             vector_index = indices.Pop();
         }
 
+        auto addrspace = var->Type()->As<core::type::Pointer>()->AddressSpace();
         core::ir::Instruction* new_access = builder.Access(
-            ir.Types().ptr(core::AddressSpace::kFunction, access_type, core::Access::kReadWrite),
-            local, indices);
+            ir.Types().ptr(addrspace, access_type, core::Access::kReadWrite), var, indices);
         new_access->InsertBefore(access);
 
         core::ir::Instruction* load = nullptr;
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
index e4adae9..657488b 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
@@ -430,5 +430,449 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToFuncParam_FromDifferentBlocks) {
+    auto* arr = b.FunctionParam(ty.array<i32, 4>());
+    auto* cond = b.FunctionParam(ty.bool_());
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func = b.Function("func", ty.i32());
+    func->SetParams({arr, cond, idx_a, idx_b});
+    b.Append(func->Block(), [&] {  //
+        auto* if_ = b.If(cond);
+        b.Append(if_->True(), [&] {  //
+            b.Return(func, b.Access(ty.i32(), arr, idx_a));
+        });
+        b.Append(if_->False(), [&] {  //
+            b.Return(func, b.Access(ty.i32(), arr, idx_b));
+        });
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+  %b1 = block {
+    if %3 [t: %b2, f: %b3] {  # if_1
+      %b2 = block {  # true
+        %6:i32 = access %2, %4
+        ret %6
+      }
+      %b3 = block {  # false
+        %7:i32 = access %2, %5
+        ret %7
+      }
+    }
+    unreachable
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+  %b1 = block {
+    %6:ptr<function, array<i32, 4>, read_write> = var, %2
+    if %3 [t: %b2, f: %b3] {  # if_1
+      %b2 = block {  # true
+        %7:ptr<function, i32, read_write> = access %6, %4
+        %8:i32 = load %7
+        ret %8
+      }
+      %b3 = block {  # false
+        %9:ptr<function, i32, read_write> = access %6, %5
+        %10:i32 = load %9
+        ret %10
+      }
+    }
+    unreachable
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+       MultipleAccessesToFuncParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+    auto* arr = b.FunctionParam(ty.array(ty.array<i32, 4>(), 4));
+    auto* cond = b.FunctionParam(ty.bool_());
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func = b.Function("func", ty.i32());
+    func->SetParams({arr, cond, idx_a, idx_b});
+    b.Append(func->Block(), [&] {  //
+        auto* if_ = b.If(cond);
+        b.Append(if_->True(), [&] {  //
+            b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+        });
+        b.Append(if_->False(), [&] {  //
+            b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+        });
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+  %b1 = block {
+    if %3 [t: %b2, f: %b3] {  # if_1
+      %b2 = block {  # true
+        %6:i32 = access %2, 0u, %4
+        ret %6
+      }
+      %b3 = block {  # false
+        %7:i32 = access %2, 0u, %5
+        ret %7
+      }
+    }
+    unreachable
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+  %b1 = block {
+    %6:array<i32, 4> = access %2, 0u
+    %7:ptr<function, array<i32, 4>, read_write> = var, %6
+    if %3 [t: %b2, f: %b3] {  # if_1
+      %b2 = block {  # true
+        %8:ptr<function, i32, read_write> = access %7, %4
+        %9:i32 = load %8
+        ret %9
+      }
+      %b3 = block {  # false
+        %10:ptr<function, i32, read_write> = access %7, %5
+        %11:i32 = load %10
+        ret %11
+      }
+    }
+    unreachable
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToBlockParam_FromDifferentBlocks) {
+    auto* arr = b.BlockParam(ty.array<i32, 4>());
+    auto* cond = b.FunctionParam(ty.bool_());
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func = b.Function("func", ty.i32());
+    func->SetParams({cond, idx_a, idx_b});
+    b.Append(func->Block(), [&] {  //
+        auto* loop = b.Loop();
+        loop->Body()->SetParams({arr});
+        b.Append(loop->Body(), [&] {
+            auto* if_ = b.If(cond);
+            b.Append(if_->True(), [&] {  //
+                b.Return(func, b.Access(ty.i32(), arr, idx_a));
+            });
+            b.Append(if_->False(), [&] {  //
+                b.Return(func, b.Access(ty.i32(), arr, idx_b));
+            });
+            b.Unreachable();
+        });
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%5:array<i32, 4>) {  # body
+        if %2 [t: %b3, f: %b4] {  # if_1
+          %b3 = block {  # true
+            %6:i32 = access %5:array<i32, 4>, %3
+            ret %6
+          }
+          %b4 = block {  # false
+            %7:i32 = access %5:array<i32, 4>, %4
+            ret %7
+          }
+        }
+        unreachable
+      }
+    }
+    unreachable
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%5:array<i32, 4>) {  # body
+        %6:ptr<function, array<i32, 4>, read_write> = var, %5:array<i32, 4>
+        if %2 [t: %b3, f: %b4] {  # if_1
+          %b3 = block {  # true
+            %7:ptr<function, i32, read_write> = access %6, %3
+            %8:i32 = load %7
+            ret %8
+          }
+          %b4 = block {  # false
+            %9:ptr<function, i32, read_write> = access %6, %4
+            %10:i32 = load %9
+            ret %10
+          }
+        }
+        unreachable
+      }
+    }
+    unreachable
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+       MultipleAccessesToBlockParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+    auto* arr = b.BlockParam(ty.array(ty.array<i32, 4>(), 4));
+    auto* cond = b.FunctionParam(ty.bool_());
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func = b.Function("func", ty.i32());
+    func->SetParams({cond, idx_a, idx_b});
+    b.Append(func->Block(), [&] {  //
+        auto* loop = b.Loop();
+        loop->Body()->SetParams({arr});
+        b.Append(loop->Body(), [&] {
+            auto* if_ = b.If(cond);
+            b.Append(if_->True(), [&] {  //
+                b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+            });
+            b.Append(if_->False(), [&] {  //
+                b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+            });
+            b.Unreachable();
+        });
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%5:array<array<i32, 4>, 4>) {  # body
+        if %2 [t: %b3, f: %b4] {  # if_1
+          %b3 = block {  # true
+            %6:i32 = access %5:array<array<i32, 4>, 4>, 0u, %3
+            ret %6
+          }
+          %b4 = block {  # false
+            %7:i32 = access %5:array<array<i32, 4>, 4>, 0u, %4
+            ret %7
+          }
+        }
+        unreachable
+      }
+    }
+    unreachable
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+  %b1 = block {
+    loop [b: %b2] {  # loop_1
+      %b2 = block (%5:array<array<i32, 4>, 4>) {  # body
+        %6:array<i32, 4> = access %5:array<array<i32, 4>, 4>, 0u
+        %7:ptr<function, array<i32, 4>, read_write> = var, %6
+        if %2 [t: %b3, f: %b4] {  # if_1
+          %b3 = block {  # true
+            %8:ptr<function, i32, read_write> = access %7, %3
+            %9:i32 = load %8
+            ret %9
+          }
+          %b4 = block {  # false
+            %10:ptr<function, i32, read_write> = access %7, %4
+            %11:i32 = load %10
+            ret %11
+          }
+        }
+        unreachable
+      }
+    }
+    unreachable
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToConstant_FromDifferentFunctions) {
+    auto* arr = b.Constant(mod.constant_values.Zero(ty.array<i32, 4>()));
+
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* func_a = b.Function("func_a", ty.i32());
+    func_a->SetParams({idx_a});
+    b.Append(func_a->Block(), [&] {  //
+        b.Return(func_a, b.Access(ty.i32(), arr, idx_a));
+    });
+
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func_b = b.Function("func_b", ty.i32());
+    func_b->SetParams({idx_b});
+    b.Append(func_b->Block(), [&] {  //
+        b.Return(func_b, b.Access(ty.i32(), arr, idx_b));
+    });
+
+    auto* idx_c = b.FunctionParam(ty.i32());
+    auto* func_c = b.Function("func_c", ty.i32());
+    func_c->SetParams({idx_c});
+    b.Append(func_c->Block(), [&] {  //
+        b.Return(func_c, b.Access(ty.i32(), arr, idx_c));
+    });
+
+    auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+  %b1 = block {
+    %3:i32 = access array<i32, 4>(0i), %2
+    ret %3
+  }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+  %b2 = block {
+    %6:i32 = access array<i32, 4>(0i), %5
+    ret %6
+  }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+  %b3 = block {
+    %9:i32 = access array<i32, 4>(0i), %8
+    ret %9
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%b1 = block {  # root
+  %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+  %b2 = block {
+    %4:ptr<private, i32, read_write> = access %1, %3
+    %5:i32 = load %4
+    ret %5
+  }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+  %b3 = block {
+    %8:ptr<private, i32, read_write> = access %1, %7
+    %9:i32 = load %8
+    ret %9
+  }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+  %b4 = block {
+    %12:ptr<private, i32, read_write> = access %1, %11
+    %13:i32 = load %12
+    ret %13
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+       MultipleAccessesToConstant_FromDifferentFunctions_WithLeadingConstantIndex) {
+    auto* arr = b.Constant(mod.constant_values.Zero(ty.array(ty.array<i32, 4>(), 4)));
+
+    auto* idx_a = b.FunctionParam(ty.i32());
+    auto* func_a = b.Function("func_a", ty.i32());
+    func_a->SetParams({idx_a});
+    b.Append(func_a->Block(), [&] {  //
+        b.Return(func_a, b.Access(ty.i32(), arr, 0_u, idx_a));
+    });
+
+    auto* idx_b = b.FunctionParam(ty.i32());
+    auto* func_b = b.Function("func_b", ty.i32());
+    func_b->SetParams({idx_b});
+    b.Append(func_b->Block(), [&] {  //
+        b.Return(func_b, b.Access(ty.i32(), arr, 0_u, idx_b));
+    });
+
+    auto* idx_c = b.FunctionParam(ty.i32());
+    auto* func_c = b.Function("func_c", ty.i32());
+    func_c->SetParams({idx_c});
+    b.Append(func_c->Block(), [&] {  //
+        b.Return(func_c, b.Access(ty.i32(), arr, 0_u, idx_c));
+    });
+
+    auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+  %b1 = block {
+    %3:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %2
+    ret %3
+  }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+  %b2 = block {
+    %6:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %5
+    ret %6
+  }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+  %b3 = block {
+    %9:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %8
+    ret %9
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%b1 = block {  # root
+  %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+  %b2 = block {
+    %4:ptr<private, i32, read_write> = access %1, %3
+    %5:i32 = load %4
+    ret %5
+  }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+  %b3 = block {
+    %8:ptr<private, i32, read_write> = access %1, %7
+    %9:i32 = load %8
+    ret %9
+  }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+  %b4 = block {
+    %12:ptr<private, i32, read_write> = access %1, %11
+    %13:i32 = load %12
+    ret %13
+  }
+}
+)";
+
+    Run(VarForDynamicIndex);
+
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
index b745f56..3b41233 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_fuzz.cc
@@ -46,7 +46,9 @@
         return;
     }
 
-    auto dst = tint::wgsl::writer::IRToProgram(ir);
+    writer::ProgramOptions program_options;
+    program_options.allowed_features = AllowedFeatures::Everything();
+    auto dst = tint::wgsl::writer::IRToProgram(ir, program_options);
     if (!dst.IsValid()) {
         std::cerr << "IR:\n" << core::ir::Disassemble(ir) << std::endl;
         if (auto result = tint::wgsl::writer::Generate(dst, {}); result == Success) {
diff --git a/src/tint/lang/wgsl/ir_roundtrip_test.cc b/src/tint/lang/wgsl/ir_roundtrip_test.cc
index 5cb3f99..6573c32 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_test.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -196,6 +196,14 @@
 )");
 }
 
+TEST_F(IRToProgramRoundtripTest, SingleFunction_UnrestrictedPointerParameters) {
+    RUN_TEST(R"(
+fn f(p : ptr<uniform, i32>) -> i32 {
+  return *(p);
+}
+)");
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // Struct declaration
 ////////////////////////////////////////////////////////////////////////////////
@@ -1921,6 +1929,21 @@
 )");
 }
 
+TEST_F(IRToProgramRoundtripTest, PhonyAssign_Conversion) {
+    RUN_TEST(R"(
+fn f() {
+  let i = 42i;
+  _ = u32(i);
+}
+)",
+             R"(
+fn f() {
+  let i = 42i;
+  _ = u32(i);
+}
+)");
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // let
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index 21782ee..94f2bac 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -35,6 +35,7 @@
 #include "src/tint/lang/core/ir/validator.h"
 #include "src/tint/lang/wgsl/builtin_fn.h"
 #include "src/tint/lang/wgsl/ir/builtin_call.h"
+#include "src/tint/utils/ice/ice.h"
 
 namespace tint::wgsl::reader {
 namespace {
@@ -166,10 +167,15 @@
         CASE(kAtomicCompareExchangeWeak)
         CASE(kSubgroupBallot)
         CASE(kSubgroupBroadcast)
-        default:
-            TINT_ICE() << "unhandled builtin function: " << fn;
-            return core::BuiltinFn::kNone;
+
+        case tint::wgsl::BuiltinFn::kBitcast:               // should lower to ir::Bitcast
+        case tint::wgsl::BuiltinFn::kWorkgroupUniformLoad:  // should be handled in Lower()
+        case tint::wgsl::BuiltinFn::kTintMaterialize:
+        case tint::wgsl::BuiltinFn::kNone:
+            break;
     }
+    TINT_ICE() << "unhandled builtin function: " << fn;
+    return core::BuiltinFn::kNone;
 }
 
 }  // namespace
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 8bf72e7..8840d93 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -68,6 +68,7 @@
 #include "src/tint/lang/core/ir/user_call.h"
 #include "src/tint/lang/core/ir/validator.h"
 #include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/texel_format.h"
 #include "src/tint/lang/core/type/atomic.h"
 #include "src/tint/lang/core/type/depth_multisampled_texture.h"
 #include "src/tint/lang/core/type/depth_texture.h"
@@ -75,7 +76,9 @@
 #include "src/tint/lang/core/type/pointer.h"
 #include "src/tint/lang/core/type/reference.h"
 #include "src/tint/lang/core/type/sampler.h"
+#include "src/tint/lang/core/type/storage_texture.h"
 #include "src/tint/lang/core/type/texture.h"
+#include "src/tint/lang/core/type/type.h"
 #include "src/tint/lang/wgsl/ir/builtin_call.h"
 #include "src/tint/lang/wgsl/ir/unary.h"
 #include "src/tint/lang/wgsl/program/program_builder.h"
@@ -968,6 +971,10 @@
                 return b.ty.sampled_texture(t->dim(), el);
             },
             [&](const core::type::StorageTexture* t) {
+                if (RequiresChromiumInternalGraphite(t)) {
+                    Enable(wgsl::Extension::kChromiumInternalGraphite);
+                }
+
                 return b.ty.storage_texture(t->dim(), t->texel_format(), t->access());
             },
             [&](const core::type::Sampler* s) { return b.ty.sampler(s->kind()); },
@@ -1051,12 +1058,17 @@
         });
     }
 
-    /// Associates the IR value @p value with the AST expression @p expr.
+    /// Associates the IR value @p value with the AST expression @p expr if it is used, otherwise
+    /// creates a phony assignment with @p expr.
     void Bind(const core::ir::Value* value, const ast::Expression* expr) {
         TINT_ASSERT(value);
-        // Value will be inlined at its place of usage.
-        if (TINT_UNLIKELY(!bindings_.Add(value, InlinedValue{expr}))) {
-            TINT_ICE() << "Bind(" << value->TypeInfo().name << ") called twice for same value";
+        if (value->IsUsed()) {
+            // Value will be inlined at its place of usage.
+            if (TINT_UNLIKELY(!bindings_.Add(value, InlinedValue{expr}))) {
+                TINT_ICE() << "Bind(" << value->TypeInfo().name << ") called twice for same value";
+            }
+        } else {
+            Append(b.Assign(b.Phony(), expr));
         }
     }
 
@@ -1192,6 +1204,12 @@
                 return false;
         }
     }
+
+    /// @returns true if the storage texture type requires the kChromiumInternalGraphite extension
+    /// to be enabled.
+    bool RequiresChromiumInternalGraphite(const core::type::StorageTexture* tex) {
+        return tex->texel_format() == core::TexelFormat::kR8Unorm;
+    }
 };
 
 }  // namespace
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 10e88ca..30e6795 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -30,8 +30,12 @@
 #include <sstream>
 #include <string>
 
+#include "src/tint/lang/core/access.h"
+#include "src/tint/lang/core/address_space.h"
 #include "src/tint/lang/core/ir/disassembler.h"
+#include "src/tint/lang/core/texel_format.h"
 #include "src/tint/lang/core/type/storage_texture.h"
+#include "src/tint/lang/core/type/texture_dimension.h"
 #include "src/tint/lang/wgsl/ir/builtin_call.h"
 #include "src/tint/lang/wgsl/ir/unary.h"
 #include "src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.h"
@@ -2629,5 +2633,23 @@
 )");
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// chromium_internal_graphite
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramTest, Enable_ChromiumInternalGraphite_SubgroupBallot) {
+    b.Append(b.ir.root_block, [&] {
+        auto t = b.Var("T", ty.ref<core::AddressSpace::kHandle>(ty.Get<core::type::StorageTexture>(
+                                core::type::TextureDimension::k2d, core::TexelFormat::kR8Unorm,
+                                core::Access::kRead, ty.f32())));
+        t->SetBindingPoint(0, 0);
+    });
+
+    EXPECT_WGSL(R"(
+enable chromium_internal_graphite;
+
+@group(0) @binding(0) var T : texture_storage_2d<r8unorm, read>;
+)");
+}
+
 }  // namespace
 }  // namespace tint::wgsl::writer
diff --git a/src/tint/lang/wgsl/writer/raise/BUILD.cmake b/src/tint/lang/wgsl/writer/raise/BUILD.cmake
index d2ae1d1..ff5b855 100644
--- a/src/tint/lang/wgsl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/wgsl/writer/raise/BUILD.cmake
@@ -112,3 +112,38 @@
 tint_target_add_external_dependencies(tint_lang_wgsl_writer_raise_test test
   "gtest"
 )
+
+################################################################################
+# Target:    tint_lang_wgsl_writer_raise_fuzz
+# Kind:      fuzz
+################################################################################
+tint_add_target(tint_lang_wgsl_writer_raise_fuzz fuzz
+  lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
+  lang/wgsl/writer/raise/raise_fuzz.cc
+  lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
+  lang/wgsl/writer/raise/value_to_let_fuzz.cc
+)
+
+tint_target_add_dependencies(tint_lang_wgsl_writer_raise_fuzz fuzz
+  tint_api_common
+  tint_cmd_fuzz_ir_fuzz
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_ir
+  tint_lang_core_type
+  tint_lang_wgsl_writer_raise
+  tint_utils_bytes
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
diff --git a/src/tint/lang/wgsl/writer/raise/BUILD.gn b/src/tint/lang/wgsl/writer/raise/BUILD.gn
index 095b41e..1f5cfa8 100644
--- a/src/tint/lang/wgsl/writer/raise/BUILD.gn
+++ b/src/tint/lang/wgsl/writer/raise/BUILD.gn
@@ -112,3 +112,35 @@
     ]
   }
 }
+
+tint_fuzz_source_set("fuzz") {
+  sources = [
+    "ptr_to_ref_fuzz.cc",
+    "raise_fuzz.cc",
+    "rename_conflicts_fuzz.cc",
+    "value_to_let_fuzz.cc",
+  ]
+  deps = [
+    "${tint_src_dir}/api/common",
+    "${tint_src_dir}/cmd/fuzz/ir:fuzz",
+    "${tint_src_dir}/lang/core",
+    "${tint_src_dir}/lang/core/constant",
+    "${tint_src_dir}/lang/core/ir",
+    "${tint_src_dir}/lang/core/type",
+    "${tint_src_dir}/lang/wgsl/writer/raise",
+    "${tint_src_dir}/utils/bytes",
+    "${tint_src_dir}/utils/containers",
+    "${tint_src_dir}/utils/diagnostic",
+    "${tint_src_dir}/utils/ice",
+    "${tint_src_dir}/utils/id",
+    "${tint_src_dir}/utils/macros",
+    "${tint_src_dir}/utils/math",
+    "${tint_src_dir}/utils/memory",
+    "${tint_src_dir}/utils/reflection",
+    "${tint_src_dir}/utils/result",
+    "${tint_src_dir}/utils/rtti",
+    "${tint_src_dir}/utils/symbol",
+    "${tint_src_dir}/utils/text",
+    "${tint_src_dir}/utils/traits",
+  ]
+}
diff --git a/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
new file mode 100644
index 0000000..b67fb22
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/ptr_to_ref_fuzz.cc
@@ -0,0 +1,50 @@
+// 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/wgsl/writer/raise/ptr_to_ref.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void PtrToRefFuzzer(core::ir::Module& module) {
+    if (auto res = PtrToRef(module); res != Success) {
+        return;
+    }
+
+    core::ir::Capabilities capabilities{core::ir::Capability::kAllowRefTypes};
+    if (auto res = Validate(module, capabilities); res != Success) {
+        TINT_ICE() << "result of PtrToRef failed IR validation\n" << res.Failure();
+    }
+}
+
+}  // namespace
+}  // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::PtrToRefFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/raise.cc b/src/tint/lang/wgsl/writer/raise/raise.cc
index 46cfe7e..dcb3e4e 100644
--- a/src/tint/lang/wgsl/writer/raise/raise.cc
+++ b/src/tint/lang/wgsl/writer/raise/raise.cc
@@ -108,6 +108,10 @@
         CASE(kPack2X16Unorm)
         CASE(kPack4X8Snorm)
         CASE(kPack4X8Unorm)
+        CASE(kPack4XI8)
+        CASE(kPack4XU8)
+        CASE(kPack4XI8Clamp)
+        CASE(kPack4XU8Clamp)
         CASE(kPow)
         CASE(kQuantizeToF16)
         CASE(kRadians)
@@ -133,6 +137,8 @@
         CASE(kUnpack2X16Unorm)
         CASE(kUnpack4X8Snorm)
         CASE(kUnpack4X8Unorm)
+        CASE(kUnpack4XI8)
+        CASE(kUnpack4XU8)
         CASE(kWorkgroupBarrier)
         CASE(kTextureBarrier)
         CASE(kTextureDimensions)
@@ -163,10 +169,11 @@
         CASE(kAtomicCompareExchangeWeak)
         CASE(kSubgroupBallot)
         CASE(kSubgroupBroadcast)
-        default:
-            TINT_ICE() << "unhandled builtin function: " << fn;
-            return wgsl::BuiltinFn::kNone;
+        case core::BuiltinFn::kNone:
+            break;
     }
+    TINT_ICE() << "unhandled builtin function: " << fn;
+    return wgsl::BuiltinFn::kNone;
 }
 
 void ReplaceBuiltinFnCall(core::ir::Module& mod, core::ir::CoreBuiltinCall* call) {
diff --git a/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc b/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc
new file mode 100644
index 0000000..e63783c
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/raise_fuzz.cc
@@ -0,0 +1,50 @@
+// 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/wgsl/writer/raise/raise.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void RaiseFuzzer(core::ir::Module& module) {
+    if (auto res = Raise(module); res != Success) {
+        return;
+    }
+
+    core::ir::Capabilities capabilities{core::ir::Capability::kAllowRefTypes};
+    if (auto res = Validate(module, capabilities); res != Success) {
+        TINT_ICE() << "result of Raise failed IR validation\n" << res.Failure();
+    }
+}
+
+}  // namespace
+}  // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::RaiseFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc b/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
new file mode 100644
index 0000000..8ac74e0
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/rename_conflicts_fuzz.cc
@@ -0,0 +1,50 @@
+// 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/wgsl/writer/raise/rename_conflicts.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void RenameConflictsFuzzer(core::ir::Module& module) {
+    if (auto res = RenameConflicts(module); res != Success) {
+        return;
+    }
+
+    core::ir::Capabilities capabilities;
+    if (auto res = Validate(module, capabilities); res != Success) {
+        TINT_ICE() << "result of RenameConflicts failed IR validation\n" << res.Failure();
+    }
+}
+
+}  // namespace
+}  // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::RenameConflictsFuzzer);
diff --git a/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc b/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc
new file mode 100644
index 0000000..84755d0
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/raise/value_to_let_fuzz.cc
@@ -0,0 +1,50 @@
+// 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/wgsl/writer/raise/value_to_let.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::wgsl::writer::raise {
+namespace {
+
+void ValueToLetFuzzer(core::ir::Module& module) {
+    if (auto res = ValueToLet(module); res != Success) {
+        return;
+    }
+
+    core::ir::Capabilities capabilities;
+    if (auto res = Validate(module, capabilities); res != Success) {
+        TINT_ICE() << "result of ValueToLet failed IR validation\n" << res.Failure();
+    }
+}
+
+}  // namespace
+}  // namespace tint::wgsl::writer::raise
+
+TINT_IR_MODULE_FUZZER(tint::wgsl::writer::raise::ValueToLetFuzzer);
diff --git a/src/tint/utils/bytes/BUILD.bazel b/src/tint/utils/bytes/BUILD.bazel
index 3b0bb76..3daeace 100644
--- a/src/tint/utils/bytes/BUILD.bazel
+++ b/src/tint/utils/bytes/BUILD.bazel
@@ -39,11 +39,13 @@
 cc_library(
   name = "bytes",
   srcs = [
+    "buffer_reader.cc",
     "bytes.cc",
     "reader.cc",
     "writer.cc",
   ],
   hdrs = [
+    "buffer_reader.h",
     "buffer_writer.h",
     "decoder.h",
     "endianness.h",
@@ -71,9 +73,9 @@
   name = "test",
   alwayslink = True,
   srcs = [
+    "buffer_reader_test.cc",
     "buffer_writer_test.cc",
     "decoder_test.cc",
-    "reader_test.cc",
     "swap_test.cc",
   ],
   deps = [
diff --git a/src/tint/utils/bytes/BUILD.cmake b/src/tint/utils/bytes/BUILD.cmake
index 6d919d3..3438bbb 100644
--- a/src/tint/utils/bytes/BUILD.cmake
+++ b/src/tint/utils/bytes/BUILD.cmake
@@ -39,6 +39,8 @@
 # Kind:      lib
 ################################################################################
 tint_add_target(tint_utils_bytes lib
+  utils/bytes/buffer_reader.cc
+  utils/bytes/buffer_reader.h
   utils/bytes/buffer_writer.h
   utils/bytes/bytes.cc
   utils/bytes/decoder.h
@@ -69,9 +71,9 @@
 # Kind:      test
 ################################################################################
 tint_add_target(tint_utils_bytes_test test
+  utils/bytes/buffer_reader_test.cc
   utils/bytes/buffer_writer_test.cc
   utils/bytes/decoder_test.cc
-  utils/bytes/reader_test.cc
   utils/bytes/swap_test.cc
 )
 
diff --git a/src/tint/utils/bytes/BUILD.gn b/src/tint/utils/bytes/BUILD.gn
index d3babe2..6618381 100644
--- a/src/tint/utils/bytes/BUILD.gn
+++ b/src/tint/utils/bytes/BUILD.gn
@@ -44,6 +44,8 @@
 
 libtint_source_set("bytes") {
   sources = [
+    "buffer_reader.cc",
+    "buffer_reader.h",
     "buffer_writer.h",
     "bytes.cc",
     "decoder.h",
@@ -71,9 +73,9 @@
 if (tint_build_unittests) {
   tint_unittests_source_set("unittests") {
     sources = [
+      "buffer_reader_test.cc",
       "buffer_writer_test.cc",
       "decoder_test.cc",
-      "reader_test.cc",
       "swap_test.cc",
     ]
     deps = [
diff --git a/src/tint/utils/bytes/buffer_reader.cc b/src/tint/utils/bytes/buffer_reader.cc
new file mode 100644
index 0000000..4294fef
--- /dev/null
+++ b/src/tint/utils/bytes/buffer_reader.cc
@@ -0,0 +1,46 @@
+// 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/utils/bytes/buffer_reader.h"
+
+namespace tint::bytes {
+
+BufferReader::~BufferReader() = default;
+
+size_t BufferReader::Read(std::byte* out, size_t count) {
+    size_t n = std::min(count, bytes_remaining_);
+    memcpy(out, data_, n);
+    data_ += n;
+    bytes_remaining_ -= n;
+    return n;
+}
+
+bool BufferReader::IsEOF() const {
+    return bytes_remaining_ == 0;
+}
+
+}  // namespace tint::bytes
diff --git a/src/tint/utils/bytes/buffer_reader.h b/src/tint/utils/bytes/buffer_reader.h
new file mode 100644
index 0000000..0367bb0
--- /dev/null
+++ b/src/tint/utils/bytes/buffer_reader.h
@@ -0,0 +1,81 @@
+// 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.
+
+#ifndef SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
+#define SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
+
+#include <algorithm>
+#include <string>
+
+#include "src/tint/utils/bytes/reader.h"
+#include "src/tint/utils/ice/ice.h"
+
+namespace tint::bytes {
+
+/// BufferReader is an implementation of the Reader interface backed by a buffer.
+class BufferReader final : public Reader {
+  public:
+    // Destructor
+    ~BufferReader() override;
+
+    /// Constructor
+    /// @param data the data to read from
+    /// @param size the number of bytes in the buffer
+    BufferReader(const std::byte* data, size_t size) : data_(data), bytes_remaining_(size) {
+        TINT_ASSERT(data);
+    }
+
+    /// Constructor
+    /// @param string the string to read from
+    explicit BufferReader(std::string_view string)
+        : data_(reinterpret_cast<const std::byte*>(string.data())),
+          bytes_remaining_(string.length()) {}
+
+    /// Constructor
+    /// @param slice the byte slice to read from
+    explicit BufferReader(Slice<const std::byte> slice)
+        : data_(slice.data), bytes_remaining_(slice.len) {
+        TINT_ASSERT(slice.data);
+    }
+
+    /// @copydoc Reader::Read
+    size_t Read(std::byte* out, size_t count) override;
+
+    /// @copydoc Reader::IsEOF
+    bool IsEOF() const override;
+
+  private:
+    /// The data to read from
+    const std::byte* data_ = nullptr;
+
+    /// The number of bytes remaining
+    size_t bytes_remaining_ = 0;
+};
+
+}  // namespace tint::bytes
+
+#endif  // SRC_TINT_UTILS_BYTES_BUFFER_READER_H_
diff --git a/src/tint/utils/bytes/reader_test.cc b/src/tint/utils/bytes/buffer_reader_test.cc
similarity index 97%
rename from src/tint/utils/bytes/reader_test.cc
rename to src/tint/utils/bytes/buffer_reader_test.cc
index 2e90e88..b0b72d8 100644
--- a/src/tint/utils/bytes/reader_test.cc
+++ b/src/tint/utils/bytes/buffer_reader_test.cc
@@ -1,4 +1,4 @@
-// Copyright 2023 The Dawn & Tint Authors
+// 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:
@@ -25,7 +25,7 @@
 // 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/utils/bytes/reader.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
 
 #include "gtest/gtest.h"
 
@@ -33,7 +33,7 @@
 namespace {
 
 template <typename... ARGS>
-auto Data(ARGS&&... args) {
+auto Data(ARGS... args) {
     return std::array{std::byte{static_cast<uint8_t>(args)}...};
 }
 
diff --git a/src/tint/utils/bytes/decoder_test.cc b/src/tint/utils/bytes/decoder_test.cc
index 125d7e6..11f663a 100644
--- a/src/tint/utils/bytes/decoder_test.cc
+++ b/src/tint/utils/bytes/decoder_test.cc
@@ -34,6 +34,7 @@
 #include <utility>
 
 #include "gmock/gmock.h"
+#include "src/tint/utils/bytes/buffer_reader.h"
 #include "src/tint/utils/result/result.h"
 
 namespace tint {
diff --git a/src/tint/utils/bytes/reader.cc b/src/tint/utils/bytes/reader.cc
index fc2dfc2..e12a020 100644
--- a/src/tint/utils/bytes/reader.cc
+++ b/src/tint/utils/bytes/reader.cc
@@ -31,18 +31,4 @@
 
 Reader::~Reader() = default;
 
-BufferReader::~BufferReader() = default;
-
-size_t BufferReader::Read(std::byte* out, size_t count) {
-    size_t n = std::min(count, bytes_remaining_);
-    memcpy(out, data_, n);
-    data_ += n;
-    bytes_remaining_ -= n;
-    return n;
-}
-
-bool BufferReader::IsEOF() const {
-    return bytes_remaining_ == 0;
-}
-
 }  // namespace tint::bytes
diff --git a/src/tint/utils/bytes/reader.h b/src/tint/utils/bytes/reader.h
index 3b5f51a..f0d074a 100644
--- a/src/tint/utils/bytes/reader.h
+++ b/src/tint/utils/bytes/reader.h
@@ -28,15 +28,11 @@
 #ifndef SRC_TINT_UTILS_BYTES_READER_H_
 #define SRC_TINT_UTILS_BYTES_READER_H_
 
-#include <algorithm>
-#include <cstdint>
 #include <string>
 
 #include "src/tint/utils/bytes/endianness.h"
 #include "src/tint/utils/bytes/swap.h"
-#include "src/tint/utils/containers/slice.h"
 #include "src/tint/utils/result/result.h"
-
 namespace tint::bytes {
 
 /// A binary stream reader interface
@@ -114,46 +110,6 @@
     }
 };
 
-/// BufferReader is an implementation of the Reader interface backed by a buffer.
-class BufferReader final : public Reader {
-  public:
-    // Destructor
-    ~BufferReader() override;
-
-    /// Constructor
-    /// @param data the data to read from
-    /// @param size the number of bytes in the buffer
-    BufferReader(const std::byte* data, size_t size) : data_(data), bytes_remaining_(size) {
-        TINT_ASSERT(data);
-    }
-
-    /// Constructor
-    /// @param string the string to read from
-    explicit BufferReader(std::string_view string)
-        : data_(reinterpret_cast<const std::byte*>(string.data())),
-          bytes_remaining_(string.length()) {}
-
-    /// Constructor
-    /// @param slice the byte slice to read from
-    explicit BufferReader(Slice<const std::byte> slice)
-        : data_(slice.data), bytes_remaining_(slice.len) {
-        TINT_ASSERT(slice.data);
-    }
-
-    /// @copydoc Reader::Read
-    size_t Read(std::byte* out, size_t count) override;
-
-    /// @copydoc Reader::IsEOF
-    bool IsEOF() const override;
-
-  private:
-    /// The data to read from
-    const std::byte* data_ = nullptr;
-
-    /// The number of bytes remaining
-    size_t bytes_remaining_ = 0;
-};
-
 }  // namespace tint::bytes
 
 #endif  // SRC_TINT_UTILS_BYTES_READER_H_