diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 337b6a1..39fae60 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -581,8 +581,8 @@
     "utils/enum_set.h",
     "utils/foreach_macro.h",
     "utils/hash.h",
-    "utils/hashmap_base.h",
     "utils/hashmap.h",
+    "utils/hashmap_base.h",
     "utils/hashset.h",
     "utils/map.h",
     "utils/math.h",
@@ -1423,7 +1423,6 @@
       "reader/wgsl/parser_impl_multiplicative_expression_test.cc",
       "reader/wgsl/parser_impl_param_list_test.cc",
       "reader/wgsl/parser_impl_paren_expression_test.cc",
-      "reader/wgsl/parser_impl_pipeline_stage_test.cc",
       "reader/wgsl/parser_impl_primary_expression_test.cc",
       "reader/wgsl/parser_impl_relational_expression_test.cc",
       "reader/wgsl/parser_impl_reserved_keyword_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index a22d9fc..1d39243 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -752,6 +752,7 @@
   endif()
 endif()
 
+
 ################################################################################
 # Tests
 ################################################################################
@@ -1054,7 +1055,6 @@
       reader/wgsl/parser_impl_multiplicative_expression_test.cc
       reader/wgsl/parser_impl_param_list_test.cc
       reader/wgsl/parser_impl_paren_expression_test.cc
-      reader/wgsl/parser_impl_pipeline_stage_test.cc
       reader/wgsl/parser_impl_primary_expression_test.cc
       reader/wgsl/parser_impl_relational_expression_test.cc
       reader/wgsl/parser_impl_reserved_keyword_test.cc
@@ -1352,7 +1352,16 @@
   add_executable(tint_unittests ${TINT_TEST_SRCS})
   set_target_properties(${target} PROPERTIES FOLDER "Tests")
 
-  if(NOT MSVC)
+  if(MSVC)
+    # TODO(crbug.com/tint/1749): MSVC debug builds can suffer from stack
+    # overflows when resolving deeply nested expression chains or statements.
+    # Production builds neither use MSVC nor debug, so just bump the stack size
+    # for this build combination.
+    string(TOUPPER "${CMAKE_BUILD_TYPE}" build_type)
+    if ((NOT ${build_type} STREQUAL "RELEASE") AND (NOT ${build_type} STREQUAL "RELWITHDEBINFO"))
+      target_link_options(tint_unittests PRIVATE "/STACK 2097152") # 2MB, default is 1MB
+    endif()
+  else()
     target_compile_options(tint_unittests PRIVATE
       -Wno-global-constructors
       -Wno-weak-vtables
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
index 05435d0..1ec3ca4 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement.cc
@@ -19,7 +19,6 @@
 #include <vector>
 
 #include "src/tint/ast/block_statement.h"
-#include "src/tint/ast/fallthrough_statement.h"
 #include "src/tint/ast/for_loop_statement.h"
 #include "src/tint/ast/if_statement.h"
 #include "src/tint/ast/loop_statement.h"
@@ -151,24 +150,9 @@
     }
 
     if (auto* case_statement = statement_node.As<ast::CaseStatement>()) {
-        // It is not OK to delete the final case statement in a switch statement if the penultimate
-        // case statement falls through to the final case statement.
-        auto* switch_statement =
-            program.Sem().Get(case_statement)->Parent()->Declaration()->As<ast::SwitchStatement>();
-
-        if (switch_statement->body.Length() > 1 &&
-            switch_statement->body[switch_statement->body.Length() - 1] == case_statement) {
-            // There are at least two case statements, and this is the final case statement.
-            auto& penultimate_case_statement_body_statements =
-                switch_statement->body[switch_statement->body.Length() - 2]->body->statements;
-            if (penultimate_case_statement_body_statements.Length() > 0 &&
-                penultimate_case_statement_body_statements
-                    [penultimate_case_statement_body_statements.Length() - 1]
-                        ->Is<ast::FallthroughStatement>()) {
-                // The penultimate case statement falls through to the final case statement, thus
-                // the final case statement cannot be removed.
-                return false;
-            }
+        // It is not OK to delete the case statement which contains the default selector.
+        if (case_statement->ContainsDefault()) {
+            return false;
         }
     }
 
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
index 925ee80..275a19e 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/delete_statement_test.cc
@@ -22,7 +22,6 @@
 #include "src/tint/ast/assignment_statement.h"
 #include "src/tint/ast/block_statement.h"
 #include "src/tint/ast/case_statement.h"
-#include "src/tint/ast/fallthrough_statement.h"
 #include "src/tint/ast/for_loop_statement.h"
 #include "src/tint/ast/if_statement.h"
 #include "src/tint/ast/switch_statement.h"
@@ -151,10 +150,7 @@
   switch(1) {
     case 0, 1: {
     }
-    default: {
-      fallthrough;
-    }
-    case 2: {
+    case 2, default: {
     }
   }
 })";
@@ -171,20 +167,14 @@
   switch(1) {
     case 0, 1: {
     }
-    default: {
-      fallthrough;
-    }
-    case 2: {
+    case 2, default: {
     }
   }
 })";
     auto expected = R"(
 fn main() {
   switch(1) {
-    default: {
-      fallthrough;
-    }
-    case 2: {
+    case 2, default: {
     }
   }
 })";
@@ -199,43 +189,6 @@
     CheckStatementDeletionWorks(original, expected, statement_finder);
 }
 
-TEST(DeleteStatementTest, DeleteFallthroughStatement) {
-    auto original = R"(
-fn main() {
-  switch(1) {
-    case 0, 1: {
-    }
-    default: {
-      fallthrough;
-    }
-    case 2: {
-    }
-  }
-})";
-    auto expected = R"(
-fn main() {
-  switch(1) {
-    case 0, 1: {
-    }
-    default: {
-    }
-    case 2: {
-    }
-  }
-})";
-    auto statement_finder = [](const Program& program) -> const ast::Statement* {
-        return program.AST()
-            .Functions()[0]
-            ->body->statements[0]
-            ->As<ast::SwitchStatement>()
-            ->body[1]
-            ->As<ast::CaseStatement>()
-            ->body->statements[0]
-            ->As<ast::FallthroughStatement>();
-    };
-    CheckStatementDeletionWorks(original, expected, statement_finder);
-}
-
 TEST(DeleteStatementTest, DeleteElse) {
     auto original = R"(
 fn main() {
@@ -570,14 +523,11 @@
     CheckStatementDeletionNotAllowed(original, statement_finder);
 }
 
-TEST(DeleteStatementTest, DoNotDeleteCaseDueToFallthrough) {
+TEST(DeleteStatementTest, DoNotDeleteCaseDueToDefault) {
     auto original = R"(
 fn main() {
   switch(1) {
-    default: {
-      fallthrough;
-    }
-    case 2: {
+    case 2, default: {
     }
   }
 })";
@@ -586,7 +536,7 @@
             .Functions()[0]
             ->body->statements[0]
             ->As<ast::SwitchStatement>()
-            ->body[1]
+            ->body[0]
             ->As<ast::CaseStatement>();
     };
     CheckStatementDeletionNotAllowed(original, statement_finder);
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
index ac5743e..8339ac0 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/replace_identifier_test.cc
@@ -386,7 +386,7 @@
 }
 
 var<private> a: S;
-let e = 3;
+const e = 3;
 @group(1) @binding(1) var<uniform> b: S;
 fn f() {
   *&a = *&b;
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index f497405..96f6d08 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -406,8 +406,8 @@
 // https://gpuweb.github.io/gpuweb/wgsl/#builtin-functions
 @const fn abs<T: fia_fiu32_f16>(T) -> T
 @const fn abs<N: num, T: fia_fiu32_f16>(vec<N, T>) -> vec<N, T>
-@const fn acos<T: fa_f32_f16>(@test_value(0.87758256189) T) -> T
-@const fn acos<N: num, T: fa_f32_f16>(@test_value(0.87758256189) vec<N, T>) -> vec<N, T>
+@const fn acos<T: fa_f32_f16>(@test_value(0.96891242171) T) -> T
+@const fn acos<N: num, T: fa_f32_f16>(@test_value(0.96891242171) vec<N, T>) -> vec<N, T>
 fn acosh<T: f32_f16>(T) -> T
 fn acosh<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 @const fn all(bool) -> bool
@@ -439,7 +439,7 @@
 @const fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 @const fn countTrailingZeros<T: iu32>(T) -> T
 @const fn countTrailingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
-fn cross<T: f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
+@const fn cross<T: fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
 fn degrees<T: f32_f16>(T) -> T
 fn degrees<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn determinant<N: num, T: f32_f16>(mat<N, N, T>) -> T
@@ -508,10 +508,10 @@
 fn modf<N: num, T: f32_f16>(vec<N, T>) -> __modf_result_vec<N, T>
 fn normalize<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn pack2x16float(vec2<f32>) -> u32
-fn pack2x16snorm(vec2<f32>) -> u32
-fn pack2x16unorm(vec2<f32>) -> u32
-fn pack4x8snorm(vec4<f32>) -> u32
-fn pack4x8unorm(vec4<f32>) -> u32
+@const fn pack2x16snorm(vec2<f32>) -> u32
+@const fn pack2x16unorm(vec2<f32>) -> u32
+@const fn pack4x8snorm(vec4<f32>) -> u32
+@const fn pack4x8unorm(vec4<f32>) -> u32
 fn pow<T: f32_f16>(T, T) -> T
 fn pow<N: num, T: f32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
 @const fn quantizeToF16(f32) -> f32
@@ -550,10 +550,10 @@
 fn trunc<T: f32_f16>(T) -> T
 fn trunc<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn unpack2x16float(u32) -> vec2<f32>
-fn unpack2x16snorm(u32) -> vec2<f32>
-fn unpack2x16unorm(u32) -> vec2<f32>
-fn unpack4x8snorm(u32) -> vec4<f32>
-fn unpack4x8unorm(u32) -> vec4<f32>
+@const fn unpack2x16snorm(u32) -> vec2<f32>
+@const fn unpack2x16unorm(u32) -> vec2<f32>
+@const fn unpack4x8snorm(u32) -> vec4<f32>
+@const fn unpack4x8unorm(u32) -> vec4<f32>
 @stage("compute") fn workgroupBarrier()
 
 fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> u32
diff --git a/src/tint/reader/spirv/enum_converter.cc b/src/tint/reader/spirv/enum_converter.cc
index 923f250..d590507 100644
--- a/src/tint/reader/spirv/enum_converter.cc
+++ b/src/tint/reader/spirv/enum_converter.cc
@@ -20,13 +20,13 @@
 
 EnumConverter::~EnumConverter() = default;
 
-ast::PipelineStage EnumConverter::ToPipelineStage(SpvExecutionModel model) {
+ast::PipelineStage EnumConverter::ToPipelineStage(spv::ExecutionModel model) {
     switch (model) {
-        case SpvExecutionModelVertex:
+        case spv::ExecutionModel::Vertex:
             return ast::PipelineStage::kVertex;
-        case SpvExecutionModelFragment:
+        case spv::ExecutionModel::Fragment:
             return ast::PipelineStage::kFragment;
-        case SpvExecutionModelGLCompute:
+        case spv::ExecutionModel::GLCompute:
             return ast::PipelineStage::kCompute;
         default:
             break;
@@ -36,23 +36,23 @@
     return ast::PipelineStage::kNone;
 }
 
-ast::AddressSpace EnumConverter::ToAddressSpace(const SpvStorageClass sc) {
+ast::AddressSpace EnumConverter::ToAddressSpace(const spv::StorageClass sc) {
     switch (sc) {
-        case SpvStorageClassInput:
+        case spv::StorageClass::Input:
             return ast::AddressSpace::kIn;
-        case SpvStorageClassOutput:
+        case spv::StorageClass::Output:
             return ast::AddressSpace::kOut;
-        case SpvStorageClassUniform:
+        case spv::StorageClass::Uniform:
             return ast::AddressSpace::kUniform;
-        case SpvStorageClassWorkgroup:
+        case spv::StorageClass::Workgroup:
             return ast::AddressSpace::kWorkgroup;
-        case SpvStorageClassUniformConstant:
+        case spv::StorageClass::UniformConstant:
             return ast::AddressSpace::kNone;
-        case SpvStorageClassStorageBuffer:
+        case spv::StorageClass::StorageBuffer:
             return ast::AddressSpace::kStorage;
-        case SpvStorageClassPrivate:
+        case spv::StorageClass::Private:
             return ast::AddressSpace::kPrivate;
-        case SpvStorageClassFunction:
+        case spv::StorageClass::Function:
             return ast::AddressSpace::kFunction;
         default:
             break;
@@ -62,31 +62,31 @@
     return ast::AddressSpace::kUndefined;
 }
 
-ast::BuiltinValue EnumConverter::ToBuiltin(SpvBuiltIn b) {
+ast::BuiltinValue EnumConverter::ToBuiltin(spv::BuiltIn b) {
     switch (b) {
-        case SpvBuiltInPosition:
+        case spv::BuiltIn::Position:
             return ast::BuiltinValue::kPosition;
-        case SpvBuiltInVertexIndex:
+        case spv::BuiltIn::VertexIndex:
             return ast::BuiltinValue::kVertexIndex;
-        case SpvBuiltInInstanceIndex:
+        case spv::BuiltIn::InstanceIndex:
             return ast::BuiltinValue::kInstanceIndex;
-        case SpvBuiltInFrontFacing:
+        case spv::BuiltIn::FrontFacing:
             return ast::BuiltinValue::kFrontFacing;
-        case SpvBuiltInFragCoord:
+        case spv::BuiltIn::FragCoord:
             return ast::BuiltinValue::kPosition;
-        case SpvBuiltInFragDepth:
+        case spv::BuiltIn::FragDepth:
             return ast::BuiltinValue::kFragDepth;
-        case SpvBuiltInLocalInvocationId:
+        case spv::BuiltIn::LocalInvocationId:
             return ast::BuiltinValue::kLocalInvocationId;
-        case SpvBuiltInLocalInvocationIndex:
+        case spv::BuiltIn::LocalInvocationIndex:
             return ast::BuiltinValue::kLocalInvocationIndex;
-        case SpvBuiltInGlobalInvocationId:
+        case spv::BuiltIn::GlobalInvocationId:
             return ast::BuiltinValue::kGlobalInvocationId;
-        case SpvBuiltInWorkgroupId:
+        case spv::BuiltIn::WorkgroupId:
             return ast::BuiltinValue::kWorkgroupId;
-        case SpvBuiltInSampleId:
+        case spv::BuiltIn::SampleId:
             return ast::BuiltinValue::kSampleIndex;
-        case SpvBuiltInSampleMask:
+        case spv::BuiltIn::SampleMask:
             return ast::BuiltinValue::kSampleMask;
         default:
             break;
@@ -96,12 +96,12 @@
     return ast::BuiltinValue::kUndefined;
 }
 
-ast::TextureDimension EnumConverter::ToDim(SpvDim dim, bool arrayed) {
+ast::TextureDimension EnumConverter::ToDim(spv::Dim dim, bool arrayed) {
     if (arrayed) {
         switch (dim) {
-            case SpvDim2D:
+            case spv::Dim::Dim2D:
                 return ast::TextureDimension::k2dArray;
-            case SpvDimCube:
+            case spv::Dim::Cube:
                 return ast::TextureDimension::kCubeArray;
             default:
                 break;
@@ -111,13 +111,13 @@
     }
     // Assume non-arrayed
     switch (dim) {
-        case SpvDim1D:
+        case spv::Dim::Dim1D:
             return ast::TextureDimension::k1d;
-        case SpvDim2D:
+        case spv::Dim::Dim2D:
             return ast::TextureDimension::k2d;
-        case SpvDim3D:
+        case spv::Dim::Dim3D:
             return ast::TextureDimension::k3d;
-        case SpvDimCube:
+        case spv::Dim::Cube:
             return ast::TextureDimension::kCube;
         default:
             break;
@@ -126,47 +126,47 @@
     return ast::TextureDimension::kNone;
 }
 
-ast::TexelFormat EnumConverter::ToTexelFormat(SpvImageFormat fmt) {
+ast::TexelFormat EnumConverter::ToTexelFormat(spv::ImageFormat fmt) {
     switch (fmt) {
-        case SpvImageFormatUnknown:
+        case spv::ImageFormat::Unknown:
             return ast::TexelFormat::kUndefined;
 
         // 8 bit channels
-        case SpvImageFormatRgba8:
+        case spv::ImageFormat::Rgba8:
             return ast::TexelFormat::kRgba8Unorm;
-        case SpvImageFormatRgba8Snorm:
+        case spv::ImageFormat::Rgba8Snorm:
             return ast::TexelFormat::kRgba8Snorm;
-        case SpvImageFormatRgba8ui:
+        case spv::ImageFormat::Rgba8ui:
             return ast::TexelFormat::kRgba8Uint;
-        case SpvImageFormatRgba8i:
+        case spv::ImageFormat::Rgba8i:
             return ast::TexelFormat::kRgba8Sint;
 
         // 16 bit channels
-        case SpvImageFormatRgba16ui:
+        case spv::ImageFormat::Rgba16ui:
             return ast::TexelFormat::kRgba16Uint;
-        case SpvImageFormatRgba16i:
+        case spv::ImageFormat::Rgba16i:
             return ast::TexelFormat::kRgba16Sint;
-        case SpvImageFormatRgba16f:
+        case spv::ImageFormat::Rgba16f:
             return ast::TexelFormat::kRgba16Float;
 
         // 32 bit channels
-        case SpvImageFormatR32ui:
+        case spv::ImageFormat::R32ui:
             return ast::TexelFormat::kR32Uint;
-        case SpvImageFormatR32i:
+        case spv::ImageFormat::R32i:
             return ast::TexelFormat::kR32Sint;
-        case SpvImageFormatR32f:
+        case spv::ImageFormat::R32f:
             return ast::TexelFormat::kR32Float;
-        case SpvImageFormatRg32ui:
+        case spv::ImageFormat::Rg32ui:
             return ast::TexelFormat::kRg32Uint;
-        case SpvImageFormatRg32i:
+        case spv::ImageFormat::Rg32i:
             return ast::TexelFormat::kRg32Sint;
-        case SpvImageFormatRg32f:
+        case spv::ImageFormat::Rg32f:
             return ast::TexelFormat::kRg32Float;
-        case SpvImageFormatRgba32ui:
+        case spv::ImageFormat::Rgba32ui:
             return ast::TexelFormat::kRgba32Uint;
-        case SpvImageFormatRgba32i:
+        case spv::ImageFormat::Rgba32i:
             return ast::TexelFormat::kRgba32Sint;
-        case SpvImageFormatRgba32f:
+        case spv::ImageFormat::Rgba32f:
             return ast::TexelFormat::kRgba32Float;
         default:
             break;
diff --git a/src/tint/reader/spirv/enum_converter.h b/src/tint/reader/spirv/enum_converter.h
index 1ebea36..0082499 100644
--- a/src/tint/reader/spirv/enum_converter.h
+++ b/src/tint/reader/spirv/enum_converter.h
@@ -16,6 +16,7 @@
 #define SRC_TINT_READER_SPIRV_ENUM_CONVERTER_H_
 
 #include "spirv/unified1/spirv.h"
+#include "spirv/unified1/spirv.hpp11"
 #include "src/tint/ast/address_space.h"
 #include "src/tint/ast/builtin_value.h"
 #include "src/tint/ast/pipeline_stage.h"
@@ -37,32 +38,49 @@
     /// On failure, logs an error and returns kNone
     /// @param model the SPIR-V entry point execution model
     /// @returns a Tint AST pipeline stage
-    ast::PipelineStage ToPipelineStage(SpvExecutionModel model);
+    ast::PipelineStage ToPipelineStage(spv::ExecutionModel model);
 
     /// Converts a SPIR-V storage class to a Tint address space.
     /// On failure, logs an error and returns kNone
     /// @param sc the SPIR-V storage class
     /// @returns a Tint AST address space
-    ast::AddressSpace ToAddressSpace(const SpvStorageClass sc);
+    ast::AddressSpace ToAddressSpace(const spv::StorageClass sc);
 
     /// Converts a SPIR-V Builtin value a Tint Builtin.
     /// On failure, logs an error and returns kNone
     /// @param b the SPIR-V builtin
     /// @returns a Tint AST builtin
-    ast::BuiltinValue ToBuiltin(SpvBuiltIn b);
+    ast::BuiltinValue ToBuiltin(spv::BuiltIn b);
 
     /// Converts a possibly arrayed SPIR-V Dim to a Tint texture dimension.
     /// On failure, logs an error and returns kNone
     /// @param dim the SPIR-V Dim value
     /// @param arrayed true if the texture is arrayed
     /// @returns a Tint AST texture dimension
-    ast::TextureDimension ToDim(SpvDim dim, bool arrayed);
+    ast::TextureDimension ToDim(spv::Dim dim, bool arrayed);
+
+    /// Converts a possibly arrayed SPIR-V Dim to a Tint texture dimension.
+    /// On failure, logs an error and returns kNone
+    /// @param dim the SPIR-V Dim value
+    /// @param arrayed true if the texture is arrayed
+    /// @returns a Tint AST texture dimension
+    ast::TextureDimension ToDim(SpvDim dim, bool arrayed) {
+        return ToDim(static_cast<spv::Dim>(dim), arrayed);
+    }
 
     /// Converts a SPIR-V Image Format to a TexelFormat
     /// On failure, logs an error and returns kNone
     /// @param fmt the SPIR-V format
     /// @returns a Tint AST format
-    ast::TexelFormat ToTexelFormat(SpvImageFormat fmt);
+    ast::TexelFormat ToTexelFormat(spv::ImageFormat fmt);
+
+    /// Converts a SPIR-V Image Format to a TexelFormat
+    /// On failure, logs an error and returns kNone
+    /// @param fmt the SPIR-V format
+    /// @returns a Tint AST format
+    ast::TexelFormat ToTexelFormat(SpvImageFormat fmt) {
+        return ToTexelFormat(static_cast<spv::ImageFormat>(fmt));
+    }
 
   private:
     /// Registers a failure and returns a stream for log diagnostics.
diff --git a/src/tint/reader/spirv/enum_converter_test.cc b/src/tint/reader/spirv/enum_converter_test.cc
index 112dc4a..b366ebd 100644
--- a/src/tint/reader/spirv/enum_converter_test.cc
+++ b/src/tint/reader/spirv/enum_converter_test.cc
@@ -24,12 +24,12 @@
 // Pipeline stage
 
 struct PipelineStageCase {
-    SpvExecutionModel model;
+    spv::ExecutionModel model;
     bool expect_success;
     ast::PipelineStage expected;
 };
 inline std::ostream& operator<<(std::ostream& out, PipelineStageCase psc) {
-    out << "PipelineStageCase{ SpvExecutionModel:" << int(psc.model)
+    out << "PipelineStageCase{ spv::ExecutionModel:::" << int(psc.model)
         << " expect_success?:" << int(psc.expect_success) << " expected:" << int(psc.expected)
         << "}";
     return out;
@@ -65,29 +65,29 @@
 
 INSTANTIATE_TEST_SUITE_P(EnumConverterGood,
                          SpvPipelineStageTest,
-                         testing::Values(PipelineStageCase{SpvExecutionModelVertex, true,
+                         testing::Values(PipelineStageCase{spv::ExecutionModel::Vertex, true,
                                                            ast::PipelineStage::kVertex},
-                                         PipelineStageCase{SpvExecutionModelFragment, true,
+                                         PipelineStageCase{spv::ExecutionModel::Fragment, true,
                                                            ast::PipelineStage::kFragment},
-                                         PipelineStageCase{SpvExecutionModelGLCompute, true,
+                                         PipelineStageCase{spv::ExecutionModel::GLCompute, true,
                                                            ast::PipelineStage::kCompute}));
 
 INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
                          SpvPipelineStageTest,
-                         testing::Values(PipelineStageCase{static_cast<SpvExecutionModel>(9999),
+                         testing::Values(PipelineStageCase{static_cast<spv::ExecutionModel>(9999),
                                                            false, ast::PipelineStage::kNone},
-                                         PipelineStageCase{SpvExecutionModelTessellationControl,
+                                         PipelineStageCase{spv::ExecutionModel::TessellationControl,
                                                            false, ast::PipelineStage::kNone}));
 
 // Storage class
 
 struct StorageClassCase {
-    SpvStorageClass sc;
+    spv::StorageClass sc;
     bool expect_success;
     ast::AddressSpace expected;
 };
 inline std::ostream& operator<<(std::ostream& out, StorageClassCase scc) {
-    out << "StorageClassCase{ SpvStorageClass:" << int(scc.sc)
+    out << "StorageClassCase{ spv::StorageClass:::" << int(scc.sc)
         << " expect_success?:" << int(scc.expect_success) << " expected:" << int(scc.expected)
         << "}";
     return out;
@@ -125,29 +125,29 @@
     EnumConverterGood,
     SpvStorageClassTest,
     testing::Values(
-        StorageClassCase{SpvStorageClassInput, true, ast::AddressSpace::kIn},
-        StorageClassCase{SpvStorageClassOutput, true, ast::AddressSpace::kOut},
-        StorageClassCase{SpvStorageClassUniform, true, ast::AddressSpace::kUniform},
-        StorageClassCase{SpvStorageClassWorkgroup, true, ast::AddressSpace::kWorkgroup},
-        StorageClassCase{SpvStorageClassUniformConstant, true, ast::AddressSpace::kNone},
-        StorageClassCase{SpvStorageClassStorageBuffer, true, ast::AddressSpace::kStorage},
-        StorageClassCase{SpvStorageClassPrivate, true, ast::AddressSpace::kPrivate},
-        StorageClassCase{SpvStorageClassFunction, true, ast::AddressSpace::kFunction}));
+        StorageClassCase{spv::StorageClass::Input, true, ast::AddressSpace::kIn},
+        StorageClassCase{spv::StorageClass::Output, true, ast::AddressSpace::kOut},
+        StorageClassCase{spv::StorageClass::Uniform, true, ast::AddressSpace::kUniform},
+        StorageClassCase{spv::StorageClass::Workgroup, true, ast::AddressSpace::kWorkgroup},
+        StorageClassCase{spv::StorageClass::UniformConstant, true, ast::AddressSpace::kNone},
+        StorageClassCase{spv::StorageClass::StorageBuffer, true, ast::AddressSpace::kStorage},
+        StorageClassCase{spv::StorageClass::Private, true, ast::AddressSpace::kPrivate},
+        StorageClassCase{spv::StorageClass::Function, true, ast::AddressSpace::kFunction}));
 
 INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
                          SpvStorageClassTest,
-                         testing::Values(StorageClassCase{static_cast<SpvStorageClass>(9999), false,
-                                                          ast::AddressSpace::kUndefined}));
+                         testing::Values(StorageClassCase{static_cast<spv::StorageClass>(9999),
+                                                          false, ast::AddressSpace::kUndefined}));
 
 // Builtin
 
 struct BuiltinCase {
-    SpvBuiltIn builtin;
+    spv::BuiltIn builtin;
     bool expect_success;
     ast::BuiltinValue expected;
 };
 inline std::ostream& operator<<(std::ostream& out, BuiltinCase bc) {
-    out << "BuiltinCase{ SpvBuiltIn:" << int(bc.builtin)
+    out << "BuiltinCase{ spv::BuiltIn::" << int(bc.builtin)
         << " expect_success?:" << int(bc.expect_success) << " expected:" << int(bc.expected) << "}";
     return out;
 }
@@ -184,43 +184,44 @@
     EnumConverterGood_Input,
     SpvBuiltinTest,
     testing::Values(
-        BuiltinCase{SpvBuiltInPosition, true, ast::BuiltinValue::kPosition},
-        BuiltinCase{SpvBuiltInInstanceIndex, true, ast::BuiltinValue::kInstanceIndex},
-        BuiltinCase{SpvBuiltInFrontFacing, true, ast::BuiltinValue::kFrontFacing},
-        BuiltinCase{SpvBuiltInFragCoord, true, ast::BuiltinValue::kPosition},
-        BuiltinCase{SpvBuiltInLocalInvocationId, true, ast::BuiltinValue::kLocalInvocationId},
-        BuiltinCase{SpvBuiltInLocalInvocationIndex, true, ast::BuiltinValue::kLocalInvocationIndex},
-        BuiltinCase{SpvBuiltInGlobalInvocationId, true, ast::BuiltinValue::kGlobalInvocationId},
-        BuiltinCase{SpvBuiltInWorkgroupId, true, ast::BuiltinValue::kWorkgroupId},
-        BuiltinCase{SpvBuiltInSampleId, true, ast::BuiltinValue::kSampleIndex},
-        BuiltinCase{SpvBuiltInSampleMask, true, ast::BuiltinValue::kSampleMask}));
+        BuiltinCase{spv::BuiltIn::Position, true, ast::BuiltinValue::kPosition},
+        BuiltinCase{spv::BuiltIn::InstanceIndex, true, ast::BuiltinValue::kInstanceIndex},
+        BuiltinCase{spv::BuiltIn::FrontFacing, true, ast::BuiltinValue::kFrontFacing},
+        BuiltinCase{spv::BuiltIn::FragCoord, true, ast::BuiltinValue::kPosition},
+        BuiltinCase{spv::BuiltIn::LocalInvocationId, true, ast::BuiltinValue::kLocalInvocationId},
+        BuiltinCase{spv::BuiltIn::LocalInvocationIndex, true,
+                    ast::BuiltinValue::kLocalInvocationIndex},
+        BuiltinCase{spv::BuiltIn::GlobalInvocationId, true, ast::BuiltinValue::kGlobalInvocationId},
+        BuiltinCase{spv::BuiltIn::WorkgroupId, true, ast::BuiltinValue::kWorkgroupId},
+        BuiltinCase{spv::BuiltIn::SampleId, true, ast::BuiltinValue::kSampleIndex},
+        BuiltinCase{spv::BuiltIn::SampleMask, true, ast::BuiltinValue::kSampleMask}));
 
 INSTANTIATE_TEST_SUITE_P(
     EnumConverterGood_Output,
     SpvBuiltinTest,
-    testing::Values(BuiltinCase{SpvBuiltInPosition, true, ast::BuiltinValue::kPosition},
-                    BuiltinCase{SpvBuiltInFragDepth, true, ast::BuiltinValue::kFragDepth},
-                    BuiltinCase{SpvBuiltInSampleMask, true, ast::BuiltinValue::kSampleMask}));
+    testing::Values(BuiltinCase{spv::BuiltIn::Position, true, ast::BuiltinValue::kPosition},
+                    BuiltinCase{spv::BuiltIn::FragDepth, true, ast::BuiltinValue::kFragDepth},
+                    BuiltinCase{spv::BuiltIn::SampleMask, true, ast::BuiltinValue::kSampleMask}));
 
 INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
                          SpvBuiltinTest,
-                         testing::Values(BuiltinCase{static_cast<SpvBuiltIn>(9999), false,
+                         testing::Values(BuiltinCase{static_cast<spv::BuiltIn>(9999), false,
                                                      ast::BuiltinValue::kUndefined},
-                                         BuiltinCase{static_cast<SpvBuiltIn>(9999), false,
+                                         BuiltinCase{static_cast<spv::BuiltIn>(9999), false,
                                                      ast::BuiltinValue::kUndefined},
-                                         BuiltinCase{SpvBuiltInNumWorkgroups, false,
+                                         BuiltinCase{spv::BuiltIn::NumWorkgroups, false,
                                                      ast::BuiltinValue::kUndefined}));
 
 // Dim
 
 struct DimCase {
-    SpvDim dim;
+    spv::Dim dim;
     bool arrayed;
     bool expect_success;
     ast::TextureDimension expected;
 };
 inline std::ostream& operator<<(std::ostream& out, DimCase dc) {
-    out << "DimCase{ SpvDim:" << int(dc.dim) << " arrayed?:" << int(dc.arrayed)
+    out << "DimCase{ spv::Dim:::" << int(dc.dim) << " arrayed?:" << int(dc.arrayed)
         << " expect_success?:" << int(dc.expect_success) << " expected:" << int(dc.expected) << "}";
     return out;
 }
@@ -256,40 +257,41 @@
                          SpvDimTest,
                          testing::Values(
                              // Non-arrayed
-                             DimCase{SpvDim1D, false, true, ast::TextureDimension::k1d},
-                             DimCase{SpvDim2D, false, true, ast::TextureDimension::k2d},
-                             DimCase{SpvDim3D, false, true, ast::TextureDimension::k3d},
-                             DimCase{SpvDimCube, false, true, ast::TextureDimension::kCube},
+                             DimCase{spv::Dim::Dim1D, false, true, ast::TextureDimension::k1d},
+                             DimCase{spv::Dim::Dim2D, false, true, ast::TextureDimension::k2d},
+                             DimCase{spv::Dim::Dim3D, false, true, ast::TextureDimension::k3d},
+                             DimCase{spv::Dim::Cube, false, true, ast::TextureDimension::kCube},
                              // Arrayed
-                             DimCase{SpvDim2D, true, true, ast::TextureDimension::k2dArray},
-                             DimCase{SpvDimCube, true, true, ast::TextureDimension::kCubeArray}));
+                             DimCase{spv::Dim::Dim2D, true, true, ast::TextureDimension::k2dArray},
+                             DimCase{spv::Dim::Cube, true, true,
+                                     ast::TextureDimension::kCubeArray}));
 
-INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
-                         SpvDimTest,
-                         testing::Values(
-                             // Invalid SPIR-V dimensionality.
-                             DimCase{SpvDimMax, false, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimMax, true, false, ast::TextureDimension::kNone},
-                             // Vulkan non-arrayed dimensionalities not supported by WGSL.
-                             DimCase{SpvDimRect, false, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimBuffer, false, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimSubpassData, false, false, ast::TextureDimension::kNone},
-                             // Arrayed dimensionalities not supported by WGSL
-                             DimCase{SpvDim3D, true, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimRect, true, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimBuffer, true, false, ast::TextureDimension::kNone},
-                             DimCase{SpvDimSubpassData, true, false,
-                                     ast::TextureDimension::kNone}));
+INSTANTIATE_TEST_SUITE_P(
+    EnumConverterBad,
+    SpvDimTest,
+    testing::Values(
+        // Invalid SPIR-V dimensionality.
+        DimCase{spv::Dim::Max, false, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::Max, true, false, ast::TextureDimension::kNone},
+        // Vulkan non-arrayed dimensionalities not supported by WGSL.
+        DimCase{spv::Dim::Rect, false, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::Buffer, false, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::SubpassData, false, false, ast::TextureDimension::kNone},
+        // Arrayed dimensionalities not supported by WGSL
+        DimCase{spv::Dim::Dim3D, true, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::Rect, true, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::Buffer, true, false, ast::TextureDimension::kNone},
+        DimCase{spv::Dim::SubpassData, true, false, ast::TextureDimension::kNone}));
 
 // TexelFormat
 
 struct TexelFormatCase {
-    SpvImageFormat format;
+    spv::ImageFormat format;
     bool expect_success;
     ast::TexelFormat expected;
 };
 inline std::ostream& operator<<(std::ostream& out, TexelFormatCase ifc) {
-    out << "TexelFormatCase{ SpvImageFormat:" << int(ifc.format)
+    out << "TexelFormatCase{ spv::ImageFormat:::" << int(ifc.format)
         << " expect_success?:" << int(ifc.expect_success) << " expected:" << int(ifc.expected)
         << "}";
     return out;
@@ -328,52 +330,52 @@
     SpvImageFormatTest,
     testing::Values(
         // Unknown.  This is used for sampled images.
-        TexelFormatCase{SpvImageFormatUnknown, true, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Unknown, true, ast::TexelFormat::kUndefined},
         // 8 bit channels
-        TexelFormatCase{SpvImageFormatRgba8, true, ast::TexelFormat::kRgba8Unorm},
-        TexelFormatCase{SpvImageFormatRgba8Snorm, true, ast::TexelFormat::kRgba8Snorm},
-        TexelFormatCase{SpvImageFormatRgba8ui, true, ast::TexelFormat::kRgba8Uint},
-        TexelFormatCase{SpvImageFormatRgba8i, true, ast::TexelFormat::kRgba8Sint},
+        TexelFormatCase{spv::ImageFormat::Rgba8, true, ast::TexelFormat::kRgba8Unorm},
+        TexelFormatCase{spv::ImageFormat::Rgba8Snorm, true, ast::TexelFormat::kRgba8Snorm},
+        TexelFormatCase{spv::ImageFormat::Rgba8ui, true, ast::TexelFormat::kRgba8Uint},
+        TexelFormatCase{spv::ImageFormat::Rgba8i, true, ast::TexelFormat::kRgba8Sint},
         // 16 bit channels
-        TexelFormatCase{SpvImageFormatRgba16ui, true, ast::TexelFormat::kRgba16Uint},
-        TexelFormatCase{SpvImageFormatRgba16i, true, ast::TexelFormat::kRgba16Sint},
-        TexelFormatCase{SpvImageFormatRgba16f, true, ast::TexelFormat::kRgba16Float},
+        TexelFormatCase{spv::ImageFormat::Rgba16ui, true, ast::TexelFormat::kRgba16Uint},
+        TexelFormatCase{spv::ImageFormat::Rgba16i, true, ast::TexelFormat::kRgba16Sint},
+        TexelFormatCase{spv::ImageFormat::Rgba16f, true, ast::TexelFormat::kRgba16Float},
         // 32 bit channels
         // ... 1 channel
-        TexelFormatCase{SpvImageFormatR32ui, true, ast::TexelFormat::kR32Uint},
-        TexelFormatCase{SpvImageFormatR32i, true, ast::TexelFormat::kR32Sint},
-        TexelFormatCase{SpvImageFormatR32f, true, ast::TexelFormat::kR32Float},
+        TexelFormatCase{spv::ImageFormat::R32ui, true, ast::TexelFormat::kR32Uint},
+        TexelFormatCase{spv::ImageFormat::R32i, true, ast::TexelFormat::kR32Sint},
+        TexelFormatCase{spv::ImageFormat::R32f, true, ast::TexelFormat::kR32Float},
         // ... 2 channels
-        TexelFormatCase{SpvImageFormatRg32ui, true, ast::TexelFormat::kRg32Uint},
-        TexelFormatCase{SpvImageFormatRg32i, true, ast::TexelFormat::kRg32Sint},
-        TexelFormatCase{SpvImageFormatRg32f, true, ast::TexelFormat::kRg32Float},
+        TexelFormatCase{spv::ImageFormat::Rg32ui, true, ast::TexelFormat::kRg32Uint},
+        TexelFormatCase{spv::ImageFormat::Rg32i, true, ast::TexelFormat::kRg32Sint},
+        TexelFormatCase{spv::ImageFormat::Rg32f, true, ast::TexelFormat::kRg32Float},
         // ... 4 channels
-        TexelFormatCase{SpvImageFormatRgba32ui, true, ast::TexelFormat::kRgba32Uint},
-        TexelFormatCase{SpvImageFormatRgba32i, true, ast::TexelFormat::kRgba32Sint},
-        TexelFormatCase{SpvImageFormatRgba32f, true, ast::TexelFormat::kRgba32Float}));
+        TexelFormatCase{spv::ImageFormat::Rgba32ui, true, ast::TexelFormat::kRgba32Uint},
+        TexelFormatCase{spv::ImageFormat::Rgba32i, true, ast::TexelFormat::kRgba32Sint},
+        TexelFormatCase{spv::ImageFormat::Rgba32f, true, ast::TexelFormat::kRgba32Float}));
 
 INSTANTIATE_TEST_SUITE_P(
     EnumConverterBad,
     SpvImageFormatTest,
     testing::Values(
         // Scanning in order from the SPIR-V spec.
-        TexelFormatCase{SpvImageFormatRg16f, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatR11fG11fB10f, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatR16f, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRgb10A2, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg16, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg8, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatR16, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatR8, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRgba16Snorm, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg16Snorm, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg8Snorm, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg16i, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg8i, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatR8i, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRgb10a2ui, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg16ui, false, ast::TexelFormat::kUndefined},
-        TexelFormatCase{SpvImageFormatRg8ui, false, ast::TexelFormat::kUndefined}));
+        TexelFormatCase{spv::ImageFormat::Rg16f, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::R11fG11fB10f, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::R16f, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rgb10A2, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg16, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg8, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::R16, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::R8, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rgba16Snorm, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg16Snorm, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg8Snorm, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg16i, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg8i, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::R8i, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rgb10a2ui, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg16ui, false, ast::TexelFormat::kUndefined},
+        TexelFormatCase{spv::ImageFormat::Rg8ui, false, ast::TexelFormat::kUndefined}));
 
 }  // namespace
 }  // namespace tint::reader::spirv
diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc
index d3febf2..0aafdf9 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -144,24 +144,99 @@
 
 namespace tint::reader::spirv {
 
+namespace three_sided_patch_function_cc {
+// This machinery is only used while SPIRV-Tools is in transition before it fully
+// uses the C++11 header spirv.hpp11
+
+/// Typedef for pointer to member function while the API call uses
+/// SpvStorageClass for its second argument.
+typedef uint32_t (
+    spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClass)(uint32_t, SpvStorageClass);
+/// Typedef for pointer to member function while the API call uses
+/// spv::StorageClass for its second argument.
+typedef uint32_t (spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClassCpp11)(
+    uint32_t,
+    spv::StorageClass);
+
+/// @param type_manager the SPIRV-Tools optimizer's type manager
+/// @param finder a pointer to member function in the type manager that does the
+/// actual lookup
+/// @param pointee_type_id the ID of the pointee type
+/// @param sc the storage class.  SC can be SpvStorageClass or spv::StorageClass
+/// @returns the ID for a SPIR-V pointer to pointee_type_id in storage class sc
+template <typename FinderType, typename SC>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_manager,
+                           FinderType finder,
+                           uint32_t pointee_type_id,
+                           SC sc);
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClass finder,
+                           uint32_t pointee_type_id,
+                           SpvStorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, sc);
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClass finder,
+                           uint32_t pointee_type_id,
+                           spv::StorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, static_cast<SpvStorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClassCpp11 finder,
+                           uint32_t pointee_type_id,
+                           SpvStorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, static_cast<spv::StorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClassCpp11 finder,
+                           uint32_t pointee_type_id,
+                           spv::StorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, sc);
+}
+}  // namespace three_sided_patch_function_cc
+
 namespace {
 
 constexpr uint32_t kMaxVectorLen = 4;
 
+template <typename FromOpcodeType>
+spv::Op ToOpcode(FromOpcodeType oc) {
+    return static_cast<spv::Op>(oc);
+}
+
+/// @param inst a SPIR-V instruction
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction& inst) {
+    return ToOpcode(inst.opcode());
+}
+/// @param inst a SPIR-V instruction pointer
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction* inst) {
+    return ToOpcode(inst->opcode());
+}
+
 // Gets the AST unary opcode for the given SPIR-V opcode, if any
 // @param opcode SPIR-V opcode
 // @param ast_unary_op return parameter
 // @returns true if it was a unary operation
-bool GetUnaryOp(SpvOp opcode, ast::UnaryOp* ast_unary_op) {
+bool GetUnaryOp(spv::Op opcode, ast::UnaryOp* ast_unary_op) {
     switch (opcode) {
-        case SpvOpSNegate:
-        case SpvOpFNegate:
+        case spv::Op::OpSNegate:
+        case spv::Op::OpFNegate:
             *ast_unary_op = ast::UnaryOp::kNegation;
             return true;
-        case SpvOpLogicalNot:
+        case spv::Op::OpLogicalNot:
             *ast_unary_op = ast::UnaryOp::kNot;
             return true;
-        case SpvOpNot:
+        case spv::Op::OpNot:
             *ast_unary_op = ast::UnaryOp::kComplement;
             return true;
         default:
@@ -173,17 +248,17 @@
 /// Converts a SPIR-V opcode for a WGSL builtin function, if there is a
 /// direct translation. Returns nullptr otherwise.
 /// @returns the WGSL builtin function name for the given opcode, or nullptr.
-const char* GetUnaryBuiltInFunctionName(SpvOp opcode) {
+const char* GetUnaryBuiltInFunctionName(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpAny:
+        case spv::Op::OpAny:
             return "any";
-        case SpvOpAll:
+        case spv::Op::OpAll:
             return "all";
-        case SpvOpIsNan:
+        case spv::Op::OpIsNan:
             return "isNan";
-        case SpvOpIsInf:
+        case spv::Op::OpIsInf:
             return "isInf";
-        case SpvOpTranspose:
+        case spv::Op::OpTranspose:
             return "transpose";
         default:
             break;
@@ -194,63 +269,63 @@
 // Converts a SPIR-V opcode to its corresponding AST binary opcode, if any
 // @param opcode SPIR-V opcode
 // @returns the AST binary op for the given opcode, or kNone
-ast::BinaryOp ConvertBinaryOp(SpvOp opcode) {
+ast::BinaryOp ConvertBinaryOp(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpIAdd:
-        case SpvOpFAdd:
+        case spv::Op::OpIAdd:
+        case spv::Op::OpFAdd:
             return ast::BinaryOp::kAdd;
-        case SpvOpISub:
-        case SpvOpFSub:
+        case spv::Op::OpISub:
+        case spv::Op::OpFSub:
             return ast::BinaryOp::kSubtract;
-        case SpvOpIMul:
-        case SpvOpFMul:
-        case SpvOpVectorTimesScalar:
-        case SpvOpMatrixTimesScalar:
-        case SpvOpVectorTimesMatrix:
-        case SpvOpMatrixTimesVector:
-        case SpvOpMatrixTimesMatrix:
+        case spv::Op::OpIMul:
+        case spv::Op::OpFMul:
+        case spv::Op::OpVectorTimesScalar:
+        case spv::Op::OpMatrixTimesScalar:
+        case spv::Op::OpVectorTimesMatrix:
+        case spv::Op::OpMatrixTimesVector:
+        case spv::Op::OpMatrixTimesMatrix:
             return ast::BinaryOp::kMultiply;
-        case SpvOpUDiv:
-        case SpvOpSDiv:
-        case SpvOpFDiv:
+        case spv::Op::OpUDiv:
+        case spv::Op::OpSDiv:
+        case spv::Op::OpFDiv:
             return ast::BinaryOp::kDivide;
-        case SpvOpUMod:
-        case SpvOpSMod:
-        case SpvOpFRem:
+        case spv::Op::OpUMod:
+        case spv::Op::OpSMod:
+        case spv::Op::OpFRem:
             return ast::BinaryOp::kModulo;
-        case SpvOpLogicalEqual:
-        case SpvOpIEqual:
-        case SpvOpFOrdEqual:
+        case spv::Op::OpLogicalEqual:
+        case spv::Op::OpIEqual:
+        case spv::Op::OpFOrdEqual:
             return ast::BinaryOp::kEqual;
-        case SpvOpLogicalNotEqual:
-        case SpvOpINotEqual:
-        case SpvOpFOrdNotEqual:
+        case spv::Op::OpLogicalNotEqual:
+        case spv::Op::OpINotEqual:
+        case spv::Op::OpFOrdNotEqual:
             return ast::BinaryOp::kNotEqual;
-        case SpvOpBitwiseAnd:
+        case spv::Op::OpBitwiseAnd:
             return ast::BinaryOp::kAnd;
-        case SpvOpBitwiseOr:
+        case spv::Op::OpBitwiseOr:
             return ast::BinaryOp::kOr;
-        case SpvOpBitwiseXor:
+        case spv::Op::OpBitwiseXor:
             return ast::BinaryOp::kXor;
-        case SpvOpLogicalAnd:
+        case spv::Op::OpLogicalAnd:
             return ast::BinaryOp::kAnd;
-        case SpvOpLogicalOr:
+        case spv::Op::OpLogicalOr:
             return ast::BinaryOp::kOr;
-        case SpvOpUGreaterThan:
-        case SpvOpSGreaterThan:
-        case SpvOpFOrdGreaterThan:
+        case spv::Op::OpUGreaterThan:
+        case spv::Op::OpSGreaterThan:
+        case spv::Op::OpFOrdGreaterThan:
             return ast::BinaryOp::kGreaterThan;
-        case SpvOpUGreaterThanEqual:
-        case SpvOpSGreaterThanEqual:
-        case SpvOpFOrdGreaterThanEqual:
+        case spv::Op::OpUGreaterThanEqual:
+        case spv::Op::OpSGreaterThanEqual:
+        case spv::Op::OpFOrdGreaterThanEqual:
             return ast::BinaryOp::kGreaterThanEqual;
-        case SpvOpULessThan:
-        case SpvOpSLessThan:
-        case SpvOpFOrdLessThan:
+        case spv::Op::OpULessThan:
+        case spv::Op::OpSLessThan:
+        case spv::Op::OpFOrdLessThan:
             return ast::BinaryOp::kLessThan;
-        case SpvOpULessThanEqual:
-        case SpvOpSLessThanEqual:
-        case SpvOpFOrdLessThanEqual:
+        case spv::Op::OpULessThanEqual:
+        case spv::Op::OpSLessThanEqual:
+        case spv::Op::OpFOrdLessThanEqual:
             return ast::BinaryOp::kLessThanEqual;
         default:
             break;
@@ -265,19 +340,19 @@
 // Othewrise returns BinaryOp::kNone.
 // @param opcode SPIR-V opcode
 // @returns operation corresponding to negated version of the SPIR-V opcode
-ast::BinaryOp NegatedFloatCompare(SpvOp opcode) {
+ast::BinaryOp NegatedFloatCompare(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpFUnordEqual:
+        case spv::Op::OpFUnordEqual:
             return ast::BinaryOp::kNotEqual;
-        case SpvOpFUnordNotEqual:
+        case spv::Op::OpFUnordNotEqual:
             return ast::BinaryOp::kEqual;
-        case SpvOpFUnordLessThan:
+        case spv::Op::OpFUnordLessThan:
             return ast::BinaryOp::kGreaterThanEqual;
-        case SpvOpFUnordLessThanEqual:
+        case spv::Op::OpFUnordLessThanEqual:
             return ast::BinaryOp::kGreaterThan;
-        case SpvOpFUnordGreaterThan:
+        case spv::Op::OpFUnordGreaterThan:
             return ast::BinaryOp::kLessThanEqual;
-        case SpvOpFUnordGreaterThanEqual:
+        case spv::Op::OpFUnordGreaterThanEqual:
             return ast::BinaryOp::kLessThan;
         default:
             break;
@@ -443,36 +518,36 @@
 
 // Returns the WGSL standard library function builtin for the
 // given instruction, or sem::BuiltinType::kNone
-sem::BuiltinType GetBuiltin(SpvOp opcode) {
+sem::BuiltinType GetBuiltin(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpBitCount:
+        case spv::Op::OpBitCount:
             return sem::BuiltinType::kCountOneBits;
-        case SpvOpBitFieldInsert:
+        case spv::Op::OpBitFieldInsert:
             return sem::BuiltinType::kInsertBits;
-        case SpvOpBitFieldSExtract:
-        case SpvOpBitFieldUExtract:
+        case spv::Op::OpBitFieldSExtract:
+        case spv::Op::OpBitFieldUExtract:
             return sem::BuiltinType::kExtractBits;
-        case SpvOpBitReverse:
+        case spv::Op::OpBitReverse:
             return sem::BuiltinType::kReverseBits;
-        case SpvOpDot:
+        case spv::Op::OpDot:
             return sem::BuiltinType::kDot;
-        case SpvOpDPdx:
+        case spv::Op::OpDPdx:
             return sem::BuiltinType::kDpdx;
-        case SpvOpDPdy:
+        case spv::Op::OpDPdy:
             return sem::BuiltinType::kDpdy;
-        case SpvOpFwidth:
+        case spv::Op::OpFwidth:
             return sem::BuiltinType::kFwidth;
-        case SpvOpDPdxFine:
+        case spv::Op::OpDPdxFine:
             return sem::BuiltinType::kDpdxFine;
-        case SpvOpDPdyFine:
+        case spv::Op::OpDPdyFine:
             return sem::BuiltinType::kDpdyFine;
-        case SpvOpFwidthFine:
+        case spv::Op::OpFwidthFine:
             return sem::BuiltinType::kFwidthFine;
-        case SpvOpDPdxCoarse:
+        case spv::Op::OpDPdxCoarse:
             return sem::BuiltinType::kDpdxCoarse;
-        case SpvOpDPdyCoarse:
+        case spv::Op::OpDPdyCoarse:
             return sem::BuiltinType::kDpdyCoarse;
-        case SpvOpFwidthCoarse:
+        case spv::Op::OpFwidthCoarse:
             return sem::BuiltinType::kFwidthCoarse;
         default:
             break;
@@ -483,20 +558,20 @@
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an image access instruction
 // whose first input operand is an OpSampledImage value.
-bool IsSampledImageAccess(SpvOp opcode) {
+bool IsSampledImageAccess(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpImageSampleImplicitLod:
-        case SpvOpImageSampleExplicitLod:
-        case SpvOpImageSampleDrefImplicitLod:
-        case SpvOpImageSampleDrefExplicitLod:
+        case spv::Op::OpImageSampleImplicitLod:
+        case spv::Op::OpImageSampleExplicitLod:
+        case spv::Op::OpImageSampleDrefImplicitLod:
+        case spv::Op::OpImageSampleDrefExplicitLod:
         // WGSL doesn't have *Proj* texturing; spirv reader emulates it.
-        case SpvOpImageSampleProjImplicitLod:
-        case SpvOpImageSampleProjExplicitLod:
-        case SpvOpImageSampleProjDrefImplicitLod:
-        case SpvOpImageSampleProjDrefExplicitLod:
-        case SpvOpImageGather:
-        case SpvOpImageDrefGather:
-        case SpvOpImageQueryLod:
+        case spv::Op::OpImageSampleProjImplicitLod:
+        case spv::Op::OpImageSampleProjExplicitLod:
+        case spv::Op::OpImageSampleProjDrefImplicitLod:
+        case spv::Op::OpImageSampleProjDrefExplicitLod:
+        case spv::Op::OpImageGather:
+        case spv::Op::OpImageDrefGather:
+        case spv::Op::OpImageQueryLod:
             return true;
         default:
             break;
@@ -506,29 +581,29 @@
 
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an atomic operation.
-bool IsAtomicOp(SpvOp opcode) {
+bool IsAtomicOp(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpAtomicLoad:
-        case SpvOpAtomicStore:
-        case SpvOpAtomicExchange:
-        case SpvOpAtomicCompareExchange:
-        case SpvOpAtomicCompareExchangeWeak:
-        case SpvOpAtomicIIncrement:
-        case SpvOpAtomicIDecrement:
-        case SpvOpAtomicIAdd:
-        case SpvOpAtomicISub:
-        case SpvOpAtomicSMin:
-        case SpvOpAtomicUMin:
-        case SpvOpAtomicSMax:
-        case SpvOpAtomicUMax:
-        case SpvOpAtomicAnd:
-        case SpvOpAtomicOr:
-        case SpvOpAtomicXor:
-        case SpvOpAtomicFlagTestAndSet:
-        case SpvOpAtomicFlagClear:
-        case SpvOpAtomicFMinEXT:
-        case SpvOpAtomicFMaxEXT:
-        case SpvOpAtomicFAddEXT:
+        case spv::Op::OpAtomicLoad:
+        case spv::Op::OpAtomicStore:
+        case spv::Op::OpAtomicExchange:
+        case spv::Op::OpAtomicCompareExchange:
+        case spv::Op::OpAtomicCompareExchangeWeak:
+        case spv::Op::OpAtomicIIncrement:
+        case spv::Op::OpAtomicIDecrement:
+        case spv::Op::OpAtomicIAdd:
+        case spv::Op::OpAtomicISub:
+        case spv::Op::OpAtomicSMin:
+        case spv::Op::OpAtomicUMin:
+        case spv::Op::OpAtomicSMax:
+        case spv::Op::OpAtomicUMax:
+        case spv::Op::OpAtomicAnd:
+        case spv::Op::OpAtomicOr:
+        case spv::Op::OpAtomicXor:
+        case spv::Op::OpAtomicFlagTestAndSet:
+        case spv::Op::OpAtomicFlagClear:
+        case spv::Op::OpAtomicFMinEXT:
+        case spv::Op::OpAtomicFMaxEXT:
+        case spv::Op::OpAtomicFAddEXT:
             return true;
         default:
             break;
@@ -539,19 +614,19 @@
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an image sampling, gather,
 // or gather-compare operation.
-bool IsImageSamplingOrGatherOrDrefGather(SpvOp opcode) {
+bool IsImageSamplingOrGatherOrDrefGather(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpImageSampleImplicitLod:
-        case SpvOpImageSampleExplicitLod:
-        case SpvOpImageSampleDrefImplicitLod:
-        case SpvOpImageSampleDrefExplicitLod:
+        case spv::Op::OpImageSampleImplicitLod:
+        case spv::Op::OpImageSampleExplicitLod:
+        case spv::Op::OpImageSampleDrefImplicitLod:
+        case spv::Op::OpImageSampleDrefExplicitLod:
             // WGSL doesn't have *Proj* texturing; spirv reader emulates it.
-        case SpvOpImageSampleProjImplicitLod:
-        case SpvOpImageSampleProjExplicitLod:
-        case SpvOpImageSampleProjDrefImplicitLod:
-        case SpvOpImageSampleProjDrefExplicitLod:
-        case SpvOpImageGather:
-        case SpvOpImageDrefGather:
+        case spv::Op::OpImageSampleProjImplicitLod:
+        case spv::Op::OpImageSampleProjExplicitLod:
+        case spv::Op::OpImageSampleProjDrefImplicitLod:
+        case spv::Op::OpImageSampleProjDrefExplicitLod:
+        case spv::Op::OpImageGather:
+        case spv::Op::OpImageDrefGather:
             return true;
         default:
             break;
@@ -562,11 +637,11 @@
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an image access instruction
 // whose first input operand is an OpImage value.
-bool IsRawImageAccess(SpvOp opcode) {
+bool IsRawImageAccess(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpImageRead:
-        case SpvOpImageWrite:
-        case SpvOpImageFetch:
+        case spv::Op::OpImageRead:
+        case spv::Op::OpImageWrite:
+        case spv::Op::OpImageFetch:
             return true;
         default:
             break;
@@ -576,13 +651,13 @@
 
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an image query instruction
-bool IsImageQuery(SpvOp opcode) {
+bool IsImageQuery(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpImageQuerySize:
-        case SpvOpImageQuerySizeLod:
-        case SpvOpImageQueryLevels:
-        case SpvOpImageQuerySamples:
-        case SpvOpImageQueryLod:
+        case spv::Op::OpImageQuerySize:
+        case spv::Op::OpImageQuerySizeLod:
+        case spv::Op::OpImageQueryLevels:
+        case spv::Op::OpImageQuerySamples:
+        case spv::Op::OpImageQueryLod:
             return true;
         default:
             break;
@@ -653,15 +728,15 @@
         // Visit successors. We will naturally skip the continue target and merge
         // blocks.
         auto* terminator = bb->terminator();
-        auto opcode = terminator->opcode();
-        if (opcode == SpvOpBranchConditional) {
+        auto opcode = ToOpcode(terminator->opcode());
+        if (opcode == spv::Op::OpBranchConditional) {
             // Visit the false branch, then the true branch, to make them come
             // out in the natural order for an "if".
             VisitBackward(terminator->GetSingleWordInOperand(2));
             VisitBackward(terminator->GetSingleWordInOperand(1));
-        } else if (opcode == SpvOpBranch) {
+        } else if (opcode == spv::Op::OpBranch) {
             VisitBackward(terminator->GetSingleWordInOperand(0));
-        } else if (opcode == SpvOpSwitch) {
+        } else if (opcode == spv::Op::OpSwitch) {
             // TODO(dneto): Consider visiting the labels in literal-value order.
             utils::Vector<uint32_t, 32> successors;
             bb->ForEachSuccessorLabel(
@@ -1265,7 +1340,7 @@
     for (uint32_t var_id : ep_info_->inputs) {
         const auto* var = def_use_mgr_->GetDef(var_id);
         TINT_ASSERT(Reader, var != nullptr);
-        TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
+        TINT_ASSERT(Reader, opcode(var) == spv::Op::OpVariable);
         auto* store_type = GetVariableStoreType(*var);
         auto* forced_param_type = store_type;
         AttributeList param_decos;
@@ -1344,7 +1419,7 @@
             } else {
                 const auto* var = def_use_mgr_->GetDef(var_id);
                 TINT_ASSERT(Reader, var != nullptr);
-                TINT_ASSERT(Reader, var->opcode() == SpvOpVariable);
+                TINT_ASSERT(Reader, opcode(var) == spv::Op::OpVariable);
                 const Type* store_type = GetVariableStoreType(*var);
                 const Type* forced_member_type = store_type;
                 AttributeList out_decos;
@@ -1570,19 +1645,19 @@
         }
 
         if (const auto* inst = block.GetMergeInst()) {
-            auto terminator_opcode = block.terminator()->opcode();
-            switch (inst->opcode()) {
-                case SpvOpSelectionMerge:
-                    if ((terminator_opcode != SpvOpBranchConditional) &&
-                        (terminator_opcode != SpvOpSwitch)) {
+            auto terminator_opcode = opcode(block.terminator());
+            switch (opcode(inst)) {
+                case spv::Op::OpSelectionMerge:
+                    if ((terminator_opcode != spv::Op::OpBranchConditional) &&
+                        (terminator_opcode != spv::Op::OpSwitch)) {
                         return Fail() << "Selection header " << block_id
                                       << " does not end in an OpBranchConditional or "
                                          "OpSwitch instruction";
                     }
                     break;
-                case SpvOpLoopMerge:
-                    if ((terminator_opcode != SpvOpBranchConditional) &&
-                        (terminator_opcode != SpvOpBranch)) {
+                case spv::Op::OpLoopMerge:
+                    if ((terminator_opcode != spv::Op::OpBranchConditional) &&
+                        (terminator_opcode != spv::Op::OpBranch)) {
                         return Fail() << "Loop header " << block_id
                                       << " does not end in an OpBranch or "
                                          "OpBranchConditional instruction";
@@ -1612,7 +1687,7 @@
             merge_info->header_for_merge = header;
             header_info->merge_for_header = merge;
 
-            if (inst->opcode() == SpvOpLoopMerge) {
+            if (opcode(inst) == spv::Op::OpLoopMerge) {
                 if (header == entry_id) {
                     return Fail() << "Function entry block " << entry_id
                                   << " cannot be a loop header";
@@ -1867,8 +1942,8 @@
                 // From the interval rule, the selection construct consists of blocks
                 // in the block order, starting at the header, until just before the
                 // merge block.
-                const auto branch_opcode = header_info->basic_block->terminator()->opcode();
-                const auto kind = (branch_opcode == SpvOpBranchConditional)
+                const auto branch_opcode = opcode(header_info->basic_block->terminator());
+                const auto kind = (branch_opcode == spv::Op::OpBranchConditional)
                                       ? Construct::kIfSelection
                                       : Construct::kSwitchSelection;
                 top = push_construct(depth, kind, header, merge);
@@ -2485,7 +2560,7 @@
         return false;
     }
     for (auto& inst : *function_.entry()) {
-        if (inst.opcode() != SpvOpVariable) {
+        if (opcode(inst) != spv::Op::OpVariable) {
             continue;
         }
         auto* var_store_type = GetVariableStoreType(inst);
@@ -2586,8 +2661,8 @@
         Fail() << "ID " << id << " does not have a defining SPIR-V instruction";
         return {};
     }
-    switch (inst->opcode()) {
-        case SpvOpVariable: {
+    switch (opcode(inst)) {
+        case spv::Op::OpVariable: {
             // This occurs for module-scope variables.
             auto name = namer_.Name(id);
             // Construct the reference type, mapping storage class correctly.
@@ -2596,7 +2671,7 @@
             return TypedExpression{type, create<ast::IdentifierExpression>(
                                              Source{}, builder_.Symbols().Register(name))};
         }
-        case SpvOpUndef:
+        case spv::Op::OpUndef:
             // Substitute a null value for undef.
             // This case occurs when OpUndef appears at module scope, as if it were
             // a constant.
@@ -3102,11 +3177,11 @@
 
 bool FunctionEmitter::EmitNormalTerminator(const BlockInfo& block_info) {
     const auto& terminator = *(block_info.basic_block->terminator());
-    switch (terminator.opcode()) {
-        case SpvOpReturn:
+    switch (opcode(terminator)) {
+        case spv::Op::OpReturn:
             AddStatement(create<ast::ReturnStatement>(Source{}));
             return true;
-        case SpvOpReturnValue: {
+        case spv::Op::OpReturnValue: {
             auto value = MakeExpression(terminator.GetSingleWordInOperand(0));
             if (!value) {
                 return false;
@@ -3114,12 +3189,12 @@
             AddStatement(create<ast::ReturnStatement>(Source{}, value.expr));
         }
             return true;
-        case SpvOpKill:
+        case spv::Op::OpKill:
             // For now, assume SPIR-V OpKill has same semantics as WGSL discard.
             // TODO(dneto): https://github.com/gpuweb/gpuweb/issues/676
             AddStatement(create<ast::DiscardStatement>(Source{}));
             return true;
-        case SpvOpUnreachable:
+        case spv::Op::OpUnreachable:
             // Translate as if it's a return. This avoids the problem where WGSL
             // requires a return statement at the end of the function body.
             {
@@ -3133,12 +3208,12 @@
                 }
             }
             return true;
-        case SpvOpBranch: {
+        case spv::Op::OpBranch: {
             const auto dest_id = terminator.GetSingleWordInOperand(0);
             AddStatement(MakeBranch(block_info, *GetBlockInfo(dest_id)));
             return true;
         }
-        case SpvOpBranchConditional: {
+        case spv::Op::OpBranchConditional: {
             // If both destinations are the same, then do the same as we would
             // for an unconditional branch (OpBranch).
             const auto true_dest = terminator.GetSingleWordInOperand(1);
@@ -3191,7 +3266,7 @@
             }
             return true;
         }
-        case SpvOpSwitch:
+        case spv::Op::OpSwitch:
             // An OpSelectionMerge must precede an OpSwitch.  That is clarified
             // in the resolution to Khronos-internal SPIR-V issue 115.
             // A new enough version of the SPIR-V validator checks this case.
@@ -3377,8 +3452,8 @@
     const auto* terminator = bb.terminator();
     const auto* merge = bb.GetMergeInst();  // Might be nullptr
     for (auto& inst : bb) {
-        if (&inst == terminator || &inst == merge || inst.opcode() == SpvOpLabel ||
-            inst.opcode() == SpvOpVariable) {
+        if (&inst == terminator || &inst == merge || opcode(inst) == spv::Op::OpLabel ||
+            opcode(inst) == spv::Op::OpVariable) {
             continue;
         }
         if (!EmitStatement(inst)) {
@@ -3546,23 +3621,23 @@
         return false;
     }
 
-    if (IsImageQuery(inst.opcode())) {
+    if (IsImageQuery(opcode(inst))) {
         return EmitImageQuery(inst);
     }
 
-    if (IsSampledImageAccess(inst.opcode()) || IsRawImageAccess(inst.opcode())) {
+    if (IsSampledImageAccess(opcode(inst)) || IsRawImageAccess(opcode(inst))) {
         return EmitImageAccess(inst);
     }
 
-    if (IsAtomicOp(inst.opcode())) {
+    if (IsAtomicOp(opcode(inst))) {
         return EmitAtomicOp(inst);
     }
 
-    switch (inst.opcode()) {
-        case SpvOpNop:
+    switch (opcode(inst)) {
+        case spv::Op::OpNop:
             return true;
 
-        case SpvOpStore: {
+        case spv::Op::OpStore: {
             auto ptr_id = inst.GetSingleWordInOperand(0);
             const auto value_id = inst.GetSingleWordInOperand(1);
 
@@ -3633,7 +3708,7 @@
             return success();
         }
 
-        case SpvOpLoad: {
+        case spv::Op::OpLoad: {
             // Memory accesses must be issued in SPIR-V program order.
             // So represent a load by a new const definition.
             const auto ptr_id = inst.GetSingleWordInOperand(0);
@@ -3684,7 +3759,7 @@
             return EmitConstDefOrWriteToHoistedVar(inst, expr);
         }
 
-        case SpvOpCopyMemory: {
+        case spv::Op::OpCopyMemory: {
             // Generate an assignment.
             auto lhs = MakeOperand(inst, 0);
             auto rhs = MakeOperand(inst, 1);
@@ -3713,7 +3788,7 @@
             return success();
         }
 
-        case SpvOpCopyObject: {
+        case spv::Op::OpCopyObject: {
             // Arguably, OpCopyObject is purely combinatorial. On the other hand,
             // it exists to make a new name for something. So we choose to make
             // a new named constant definition.
@@ -3732,39 +3807,39 @@
             return EmitConstDefOrWriteToHoistedVar(inst, expr);
         }
 
-        case SpvOpPhi: {
+        case spv::Op::OpPhi: {
             // The value will be in scope, available for reading from the phi ID.
             return true;
         }
 
-        case SpvOpOuterProduct:
+        case spv::Op::OpOuterProduct:
             // Synthesize an outer product expression in its own statement.
             return EmitConstDefOrWriteToHoistedVar(inst, MakeOuterProduct(inst));
 
-        case SpvOpVectorInsertDynamic:
+        case spv::Op::OpVectorInsertDynamic:
             // Synthesize a vector insertion in its own statements.
             return MakeVectorInsertDynamic(inst);
 
-        case SpvOpCompositeInsert:
+        case spv::Op::OpCompositeInsert:
             // Synthesize a composite insertion in its own statements.
             return MakeCompositeInsert(inst);
 
-        case SpvOpFunctionCall:
+        case spv::Op::OpFunctionCall:
             return EmitFunctionCall(inst);
 
-        case SpvOpControlBarrier:
+        case spv::Op::OpControlBarrier:
             return EmitControlBarrier(inst);
 
-        case SpvOpExtInst:
+        case spv::Op::OpExtInst:
             if (parser_impl_.IsIgnoredExtendedInstruction(inst)) {
                 return true;
             }
             break;
 
-        case SpvOpIAddCarry:
-        case SpvOpISubBorrow:
-        case SpvOpUMulExtended:
-        case SpvOpSMulExtended:
+        case spv::Op::OpIAddCarry:
+        case spv::Op::OpISubBorrow:
+        case spv::Op::OpUMulExtended:
+        case spv::Op::OpSMulExtended:
             return Fail() << "extended arithmetic is not finalized for WGSL: "
                              "https://github.com/gpuweb/gpuweb/issues/1565: "
                           << inst.PrettyPrint();
@@ -3772,7 +3847,7 @@
         default:
             break;
     }
-    return Fail() << "unhandled instruction with opcode " << inst.opcode() << ": "
+    return Fail() << "unhandled instruction with opcode " << uint32_t(opcode(inst)) << ": "
                   << inst.PrettyPrint();
 }
 
@@ -3791,7 +3866,7 @@
         return {};
     }
 
-    const auto opcode = inst.opcode();
+    const auto op = opcode(inst);
 
     const Type* ast_type = nullptr;
     if (inst.type_id()) {
@@ -3802,7 +3877,7 @@
         }
     }
 
-    auto binary_op = ConvertBinaryOp(opcode);
+    auto binary_op = ConvertBinaryOp(op);
     if (binary_op != ast::BinaryOp::kNone) {
         auto arg0 = MakeOperand(inst, 0);
         auto arg1 =
@@ -3817,14 +3892,14 @@
     }
 
     auto unary_op = ast::UnaryOp::kNegation;
-    if (GetUnaryOp(opcode, &unary_op)) {
+    if (GetUnaryOp(op, &unary_op)) {
         auto arg0 = MakeOperand(inst, 0);
         auto* unary_expr = create<ast::UnaryOpExpression>(Source{}, unary_op, arg0.expr);
         TypedExpression result{ast_type, unary_expr};
         return parser_impl_.RectifyForcedResultType(result, inst, arg0.type);
     }
 
-    const char* unary_builtin_name = GetUnaryBuiltInFunctionName(opcode);
+    const char* unary_builtin_name = GetUnaryBuiltInFunctionName(op);
     if (unary_builtin_name != nullptr) {
         ExpressionList params;
         params.Push(MakeOperand(inst, 0).expr);
@@ -3835,40 +3910,40 @@
                               std::move(params))};
     }
 
-    const auto builtin = GetBuiltin(opcode);
+    const auto builtin = GetBuiltin(op);
     if (builtin != sem::BuiltinType::kNone) {
         return MakeBuiltinCall(inst);
     }
 
-    if (opcode == SpvOpFMod) {
+    if (op == spv::Op::OpFMod) {
         return MakeFMod(inst);
     }
 
-    if (opcode == SpvOpAccessChain || opcode == SpvOpInBoundsAccessChain) {
+    if (op == spv::Op::OpAccessChain || op == spv::Op::OpInBoundsAccessChain) {
         return MakeAccessChain(inst);
     }
 
-    if (opcode == SpvOpBitcast) {
+    if (op == spv::Op::OpBitcast) {
         return {ast_type, create<ast::BitcastExpression>(Source{}, ast_type->Build(builder_),
                                                          MakeOperand(inst, 0).expr)};
     }
 
-    if (opcode == SpvOpShiftLeftLogical || opcode == SpvOpShiftRightLogical ||
-        opcode == SpvOpShiftRightArithmetic) {
+    if (op == spv::Op::OpShiftLeftLogical || op == spv::Op::OpShiftRightLogical ||
+        op == spv::Op::OpShiftRightArithmetic) {
         auto arg0 = MakeOperand(inst, 0);
         // The second operand must be unsigned. It's ok to wrap the shift amount
         // since the shift is modulo the bit width of the first operand.
         auto arg1 = parser_impl_.AsUnsigned(MakeOperand(inst, 1));
 
-        switch (opcode) {
-            case SpvOpShiftLeftLogical:
+        switch (op) {
+            case spv::Op::OpShiftLeftLogical:
                 binary_op = ast::BinaryOp::kShiftLeft;
                 break;
-            case SpvOpShiftRightLogical:
+            case spv::Op::OpShiftRightLogical:
                 arg0 = parser_impl_.AsUnsigned(arg0);
                 binary_op = ast::BinaryOp::kShiftRight;
                 break;
-            case SpvOpShiftRightArithmetic:
+            case spv::Op::OpShiftRightArithmetic:
                 arg0 = parser_impl_.AsSigned(arg0);
                 binary_op = ast::BinaryOp::kShiftRight;
                 break;
@@ -3880,7 +3955,7 @@
         return parser_impl_.RectifyForcedResultType(result, inst, arg0.type);
     }
 
-    auto negated_op = NegatedFloatCompare(opcode);
+    auto negated_op = NegatedFloatCompare(op);
     if (negated_op != ast::BinaryOp::kNone) {
         auto arg0 = MakeOperand(inst, 0);
         auto arg1 = MakeOperand(inst, 1);
@@ -3891,7 +3966,7 @@
         return {ast_type, negated_expr};
     }
 
-    if (opcode == SpvOpExtInst) {
+    if (op == spv::Op::OpExtInst) {
         if (parser_impl_.IsIgnoredExtendedInstruction(inst)) {
             // Ignore it but don't error out.
             return {};
@@ -3904,7 +3979,7 @@
         return EmitGlslStd450ExtInst(inst);
     }
 
-    if (opcode == SpvOpCompositeConstruct) {
+    if (op == spv::Op::OpCompositeConstruct) {
         ExpressionList operands;
         for (uint32_t iarg = 0; iarg < inst.NumInOperands(); ++iarg) {
             operands.Push(MakeOperand(inst, iarg).expr);
@@ -3913,34 +3988,34 @@
                 builder_.Construct(Source{}, ast_type->Build(builder_), std::move(operands))};
     }
 
-    if (opcode == SpvOpCompositeExtract) {
+    if (op == spv::Op::OpCompositeExtract) {
         return MakeCompositeExtract(inst);
     }
 
-    if (opcode == SpvOpVectorShuffle) {
+    if (op == spv::Op::OpVectorShuffle) {
         return MakeVectorShuffle(inst);
     }
 
-    if (opcode == SpvOpVectorExtractDynamic) {
+    if (op == spv::Op::OpVectorExtractDynamic) {
         return {ast_type, create<ast::IndexAccessorExpression>(Source{}, MakeOperand(inst, 0).expr,
                                                                MakeOperand(inst, 1).expr)};
     }
 
-    if (opcode == SpvOpConvertSToF || opcode == SpvOpConvertUToF || opcode == SpvOpConvertFToS ||
-        opcode == SpvOpConvertFToU) {
+    if (op == spv::Op::OpConvertSToF || op == spv::Op::OpConvertUToF ||
+        op == spv::Op::OpConvertFToS || op == spv::Op::OpConvertFToU) {
         return MakeNumericConversion(inst);
     }
 
-    if (opcode == SpvOpUndef) {
+    if (op == spv::Op::OpUndef) {
         // Replace undef with the null value.
         return parser_impl_.MakeNullExpression(ast_type);
     }
 
-    if (opcode == SpvOpSelect) {
+    if (op == spv::Op::OpSelect) {
         return MakeSimpleSelect(inst);
     }
 
-    if (opcode == SpvOpArrayLength) {
+    if (op == spv::Op::OpArrayLength) {
         return MakeArrayLength(inst);
     }
 
@@ -4424,12 +4499,12 @@
     const auto constants = constant_mgr_->GetOperandConstants(&inst);
 
     const auto* ptr_type_inst = def_use_mgr_->GetDef(ptr_ty_id);
-    if (!ptr_type_inst || (ptr_type_inst->opcode() != SpvOpTypePointer)) {
+    if (!ptr_type_inst || (opcode(ptr_type_inst) != spv::Op::OpTypePointer)) {
         Fail() << "Access chain %" << inst.result_id() << " base pointer is not of pointer type";
         return {};
     }
-    SpvStorageClass address_space =
-        static_cast<SpvStorageClass>(ptr_type_inst->GetSingleWordInOperand(0));
+    spv::StorageClass address_space =
+        static_cast<spv::StorageClass>(ptr_type_inst->GetSingleWordInOperand(0));
     uint32_t pointee_type_id = ptr_type_inst->GetSingleWordInOperand(1);
 
     // Build up a nested expression for the access chain by walking down the type
@@ -4446,8 +4521,8 @@
                    << (index - first_index) << " indices: " << inst.PrettyPrint();
             return {};
         }
-        switch (pointee_type_inst->opcode()) {
-            case SpvOpTypeVector:
+        switch (opcode(pointee_type_inst)) {
+            case spv::Op::OpTypeVector:
                 if (index_const) {
                     // Try generating a MemberAccessor expression
                     const auto num_elems = pointee_type_inst->GetSingleWordInOperand(1);
@@ -4473,24 +4548,24 @@
                 // Sink pointers to vector components.
                 sink_pointer = true;
                 break;
-            case SpvOpTypeMatrix:
+            case spv::Op::OpTypeMatrix:
                 // Use array syntax.
                 next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
                                                                  MakeOperand(inst, index).expr);
                 // All matrix components are the same type.
                 pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpTypeArray:
+            case spv::Op::OpTypeArray:
                 next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
                                                                  MakeOperand(inst, index).expr);
                 pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpTypeRuntimeArray:
+            case spv::Op::OpTypeRuntimeArray:
                 next_expr = create<ast::IndexAccessorExpression>(Source{}, current_expr.expr,
                                                                  MakeOperand(inst, index).expr);
                 pointee_type_id = pointee_type_inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpTypeStruct: {
+            case spv::Op::OpTypeStruct: {
                 if (!index_const) {
                     Fail() << "Access chain %" << inst.result_id() << " index %"
                            << inst.GetSingleWordInOperand(index)
@@ -4519,7 +4594,9 @@
                        << ": " << pointee_type_inst->PrettyPrint();
                 return {};
         }
-        const auto pointer_type_id = type_mgr_->FindPointerToType(pointee_type_id, address_space);
+        const auto pointer_type_id = three_sided_patch_function_cc::FindPointerToType(
+            type_mgr_, &spvtools::opt::analysis::TypeManager::FindPointerToType, pointee_type_id,
+            address_space);
         auto* type = parser_impl_.ConvertType(pointer_type_id, PtrAs::Ref);
         TINT_ASSERT(Reader, type && type->Is<Reference>());
         current_expr = TypedExpression{type, next_expr};
@@ -4594,11 +4671,11 @@
             return {};
         }
         const char* operation_name = nullptr;
-        switch (inst.opcode()) {
-            case SpvOpCompositeExtract:
+        switch (opcode(inst)) {
+            case spv::Op::OpCompositeExtract:
                 operation_name = "OpCompositeExtract";
                 break;
-            case SpvOpCompositeInsert:
+            case spv::Op::OpCompositeInsert:
                 operation_name = "OpCompositeInsert";
                 break;
             default:
@@ -4606,8 +4683,8 @@
                 return {};
         }
         const ast::Expression* next_expr = nullptr;
-        switch (current_type_inst->opcode()) {
-            case SpvOpTypeVector: {
+        switch (opcode(current_type_inst)) {
+            case spv::Op::OpTypeVector: {
                 // Try generating a MemberAccessor expression. That result in something
                 // like  "foo.z", which is more idiomatic than "foo[2]".
                 const auto num_elems = current_type_inst->GetSingleWordInOperand(1);
@@ -4628,7 +4705,7 @@
                 current_type_id = current_type_inst->GetSingleWordInOperand(0);
                 break;
             }
-            case SpvOpTypeMatrix: {
+            case spv::Op::OpTypeMatrix: {
                 // Check bounds
                 const auto num_elems = current_type_inst->GetSingleWordInOperand(1);
                 if (num_elems <= index_val) {
@@ -4648,7 +4725,7 @@
                 current_type_id = current_type_inst->GetSingleWordInOperand(0);
                 break;
             }
-            case SpvOpTypeArray:
+            case spv::Op::OpTypeArray:
                 // The array size could be a spec constant, and so it's not always
                 // statically checkable.  Instead, rely on a runtime index clamp
                 // or runtime check to keep this safe.
@@ -4656,11 +4733,11 @@
                                                                  make_index(index_val));
                 current_type_id = current_type_inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpTypeRuntimeArray:
+            case spv::Op::OpTypeRuntimeArray:
                 Fail() << "can't do " << operation_name
                        << " on a runtime array: " << inst.PrettyPrint();
                 return {};
-            case SpvOpTypeStruct: {
+            case spv::Op::OpTypeStruct: {
                 const auto num_members = current_type_inst->NumInOperands();
                 if (num_members <= index_val) {
                     Fail() << operation_name << " %" << inst.result_id() << " index value "
@@ -4751,14 +4828,14 @@
         auto& def = def_info_[id];
         // Builtins are always defined outside the function.
         switch (builtin) {
-            case SpvBuiltInPointSize:
+            case spv::BuiltIn::PointSize:
                 def->skip = SkipReason::kPointSizeBuiltinPointer;
                 break;
-            case SpvBuiltInSampleMask: {
+            case spv::BuiltIn::SampleMask: {
                 // Distinguish between input and output variable.
                 const auto storage_class =
-                    static_cast<SpvStorageClass>(var->GetSingleWordInOperand(0));
-                if (storage_class == SpvStorageClassInput) {
+                    static_cast<spv::StorageClass>(var->GetSingleWordInOperand(0));
+                if (storage_class == spv::StorageClass::Input) {
                     sample_mask_in_id = id;
                     def->skip = SkipReason::kSampleMaskInBuiltinPointer;
                 } else {
@@ -4767,14 +4844,14 @@
                 }
                 break;
             }
-            case SpvBuiltInSampleId:
-            case SpvBuiltInInstanceIndex:
-            case SpvBuiltInVertexIndex:
-            case SpvBuiltInLocalInvocationIndex:
-            case SpvBuiltInLocalInvocationId:
-            case SpvBuiltInGlobalInvocationId:
-            case SpvBuiltInWorkgroupId:
-            case SpvBuiltInNumWorkgroups:
+            case spv::BuiltIn::SampleId:
+            case spv::BuiltIn::InstanceIndex:
+            case spv::BuiltIn::VertexIndex:
+            case spv::BuiltIn::LocalInvocationIndex:
+            case spv::BuiltIn::LocalInvocationId:
+            case spv::BuiltIn::GlobalInvocationId:
+            case spv::BuiltIn::WorkgroupId:
+            case spv::BuiltIn::NumWorkgroups:
                 break;
             default:
                 return Fail() << "unrecognized special builtin: " << int(builtin);
@@ -4791,7 +4868,7 @@
         const auto block_pos = block_info->pos;
         for (const auto& inst : *(block_info->basic_block)) {
             const auto result_id = inst.result_id();
-            if ((result_id == 0) || inst.opcode() == SpvOpLabel) {
+            if ((result_id == 0) || opcode(inst) == spv::Op::OpLabel) {
                 continue;
             }
             def_info_[result_id] = std::make_unique<DefInfo>(index, inst, block_pos);
@@ -4807,15 +4884,15 @@
                 // pointer definitions are SSA values, and their definitions must be
                 // visited before their uses.
                 if (type->AsPointer()) {
-                    switch (inst.opcode()) {
-                        case SpvOpUndef:
+                    switch (opcode(inst)) {
+                        case spv::Op::OpUndef:
                             return Fail() << "undef pointer is not valid: " << inst.PrettyPrint();
-                        case SpvOpVariable:
+                        case spv::Op::OpVariable:
                             info->pointer = GetPointerInfo(result_id);
                             break;
-                        case SpvOpAccessChain:
-                        case SpvOpInBoundsAccessChain:
-                        case SpvOpCopyObject:
+                        case spv::Op::OpAccessChain:
+                        case spv::Op::OpInBoundsAccessChain:
+                        case spv::Op::OpCopyObject:
                             // Inherit from the first operand. We need this so we can pick up
                             // a remapped storage buffer.
                             info->pointer = GetPointerInfo(inst.GetSingleWordInOperand(0));
@@ -4846,8 +4923,8 @@
         [&](const spvtools::opt::Instruction& inst) -> DefInfo::Pointer {
         // WGSL root identifiers (or SPIR-V "memory object declarations") are
         // either variables or function parameters.
-        switch (inst.opcode()) {
-            case SpvOpVariable: {
+        switch (opcode(inst)) {
+            case spv::Op::OpVariable: {
                 if (const auto* module_var = parser_impl_.GetModuleVariable(id)) {
                     return DefInfo::Pointer{module_var->declared_address_space,
                                             module_var->declared_access};
@@ -4856,7 +4933,7 @@
                 // access mode.
                 return DefInfo::Pointer{ast::AddressSpace::kFunction, ast::Access::kUndefined};
             }
-            case SpvOpFunctionParameter: {
+            case spv::Op::OpFunctionParameter: {
                 const auto* type = As<Pointer>(parser_impl_.ConvertType(inst.type_id()));
                 // For access mode, kUndefined is ok for now, since the
                 // only non-default access mode on a pointer would be for a storage
@@ -4879,7 +4956,7 @@
     auto where = def_info_.find(id);
     if (where != def_info_.end()) {
         const auto& info = where->second;
-        if (info->inst.opcode() == SpvOpVariable) {
+        if (opcode(info->inst) == spv::Op::OpVariable) {
             // Ignore the cache in this case and compute it from scratch.
             // That's because for a function-scope OpVariable is a
             // locally-defined value.  So its cache entry has been created
@@ -4925,8 +5002,8 @@
     };
     for (auto& id_def_info_pair : def_info_) {
         const auto& inst = id_def_info_pair.second->inst;
-        const auto opcode = inst.opcode();
-        if ((opcode == SpvOpVectorShuffle) || (opcode == SpvOpOuterProduct)) {
+        const auto op = opcode(inst);
+        if ((op == spv::Op::OpVectorShuffle) || (op == spv::Op::OpOuterProduct)) {
             // We might access the vector operands multiple times. Make sure they
             // are evaluated only once.
             require_named_const_def(inst, 0);
@@ -4979,7 +5056,7 @@
         const auto* block_info = GetBlockInfo(block_id);
         for (const auto& inst : *(block_info->basic_block)) {
             // Update bookkeeping for locally-defined IDs used by this instruction.
-            if (inst.opcode() == SpvOpPhi) {
+            if (opcode(inst) == spv::Op::OpPhi) {
                 // For an OpPhi defining value P, an incoming value V from parent block B is
                 // counted as being "used" at block B, not at the block containing the Phi.
                 // That's because we will create a variable PHI_P to hold the phi value, and
@@ -5128,7 +5205,7 @@
 }
 
 TypedExpression FunctionEmitter::MakeNumericConversion(const spvtools::opt::Instruction& inst) {
-    const auto opcode = inst.opcode();
+    const auto op = opcode(inst);
     auto* requested_type = parser_impl_.ConvertType(inst.type_id());
     auto arg_expr = MakeOperand(inst, 0);
     if (!arg_expr) {
@@ -5137,7 +5214,7 @@
     arg_expr.type = arg_expr.type->UnwrapRef();
 
     const Type* expr_type = nullptr;
-    if ((opcode == SpvOpConvertSToF) || (opcode == SpvOpConvertUToF)) {
+    if ((op == spv::Op::OpConvertSToF) || (op == spv::Op::OpConvertUToF)) {
         if (arg_expr.type->IsIntegerScalarOrVector()) {
             expr_type = requested_type;
         } else {
@@ -5145,7 +5222,7 @@
                       "scalar or vector: "
                    << inst.PrettyPrint();
         }
-    } else if (inst.opcode() == SpvOpConvertFToU) {
+    } else if (op == spv::Op::OpConvertFToU) {
         if (arg_expr.type->IsFloatScalarOrVector()) {
             expr_type = parser_impl_.GetUnsignedIntMatchingShape(arg_expr.type);
         } else {
@@ -5153,7 +5230,7 @@
                       "point scalar or vector: "
                    << inst.PrettyPrint();
         }
-    } else if (inst.opcode() == SpvOpConvertFToS) {
+    } else if (op == spv::Op::OpConvertFToS) {
         if (arg_expr.type->IsFloatScalarOrVector()) {
             expr_type = parser_impl_.GetSignedIntMatchingShape(arg_expr.type);
         } else {
@@ -5228,28 +5305,28 @@
     uint32_t memory = operands[1];
     uint32_t semantics = operands[2];
 
-    if (execution != SpvScopeWorkgroup) {
+    if (execution != uint32_t(spv::Scope::Workgroup)) {
         return Fail() << "unsupported control barrier execution scope: "
                       << "expected Workgroup (2), got: " << execution;
     }
-    if (semantics & SpvMemorySemanticsAcquireReleaseMask) {
-        semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsAcquireReleaseMask);
+    if (semantics & uint32_t(spv::MemorySemanticsMask::AcquireRelease)) {
+        semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::AcquireRelease);
     } else {
         return Fail() << "control barrier semantics requires acquire and release";
     }
-    if (semantics & SpvMemorySemanticsWorkgroupMemoryMask) {
-        if (memory != SpvScopeWorkgroup) {
+    if (semantics & uint32_t(spv::MemorySemanticsMask::WorkgroupMemory)) {
+        if (memory != uint32_t(spv::Scope::Workgroup)) {
             return Fail() << "workgroupBarrier requires workgroup memory scope";
         }
         AddStatement(create<ast::CallStatement>(builder_.Call("workgroupBarrier")));
-        semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsWorkgroupMemoryMask);
+        semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::WorkgroupMemory);
     }
-    if (semantics & SpvMemorySemanticsUniformMemoryMask) {
-        if (memory != SpvScopeDevice) {
+    if (semantics & uint32_t(spv::MemorySemanticsMask::UniformMemory)) {
+        if (memory != uint32_t(spv::Scope::Device)) {
             return Fail() << "storageBarrier requires device memory scope";
         }
         AddStatement(create<ast::CallStatement>(builder_.Call("storageBarrier")));
-        semantics &= ~static_cast<uint32_t>(SpvMemorySemanticsUniformMemoryMask);
+        semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::UniformMemory);
     }
     if (semantics) {
         return Fail() << "unsupported control barrier semantics: " << semantics;
@@ -5258,7 +5335,7 @@
 }
 
 TypedExpression FunctionEmitter::MakeBuiltinCall(const spvtools::opt::Instruction& inst) {
-    const auto builtin = GetBuiltin(inst.opcode());
+    const auto builtin = GetBuiltin(opcode(inst));
     auto* name = sem::str(builtin);
     auto* ident = create<ast::IdentifierExpression>(Source{}, builder_.Symbols().Register(name));
 
@@ -5374,7 +5451,7 @@
 
 bool FunctionEmitter::EmitImageAccess(const spvtools::opt::Instruction& inst) {
     ExpressionList args;
-    const auto opcode = inst.opcode();
+    const auto op = opcode(inst);
 
     // Form the texture operand.
     const spvtools::opt::Instruction* image = GetImage(inst);
@@ -5384,7 +5461,7 @@
     args.Push(GetImageExpression(inst));
 
     // Form the sampler operand, if needed.
-    if (IsSampledImageAccess(opcode)) {
+    if (IsSampledImageAccess(op)) {
         // Form the sampler operand.
         if (auto* sampler = GetSamplerExpression(inst)) {
             args.Push(sampler);
@@ -5438,25 +5515,25 @@
     bool is_dref_sample = false;
     bool is_gather_or_dref_gather = false;
     bool is_non_dref_sample = false;
-    switch (opcode) {
-        case SpvOpImageSampleImplicitLod:
-        case SpvOpImageSampleExplicitLod:
-        case SpvOpImageSampleProjImplicitLod:
-        case SpvOpImageSampleProjExplicitLod:
+    switch (op) {
+        case spv::Op::OpImageSampleImplicitLod:
+        case spv::Op::OpImageSampleExplicitLod:
+        case spv::Op::OpImageSampleProjImplicitLod:
+        case spv::Op::OpImageSampleProjExplicitLod:
             is_non_dref_sample = true;
             builtin_name = "textureSample";
             break;
-        case SpvOpImageSampleDrefImplicitLod:
-        case SpvOpImageSampleDrefExplicitLod:
-        case SpvOpImageSampleProjDrefImplicitLod:
-        case SpvOpImageSampleProjDrefExplicitLod:
+        case spv::Op::OpImageSampleDrefImplicitLod:
+        case spv::Op::OpImageSampleDrefExplicitLod:
+        case spv::Op::OpImageSampleProjDrefImplicitLod:
+        case spv::Op::OpImageSampleProjDrefExplicitLod:
             is_dref_sample = true;
             builtin_name = "textureSampleCompare";
             if (!consume_dref()) {
                 return false;
             }
             break;
-        case SpvOpImageGather:
+        case spv::Op::OpImageGather:
             is_gather_or_dref_gather = true;
             builtin_name = "textureGather";
             if (!texture_type->Is<DepthTexture>()) {
@@ -5471,20 +5548,20 @@
             // Skip over the component operand, even for depth textures.
             arg_index++;
             break;
-        case SpvOpImageDrefGather:
+        case spv::Op::OpImageDrefGather:
             is_gather_or_dref_gather = true;
             builtin_name = "textureGatherCompare";
             if (!consume_dref()) {
                 return false;
             }
             break;
-        case SpvOpImageFetch:
-        case SpvOpImageRead:
+        case spv::Op::OpImageFetch:
+        case spv::Op::OpImageRead:
             // Read a single texel from a sampled or storage image.
             builtin_name = "textureLoad";
             use_level_of_detail_suffix = false;
             break;
-        case SpvOpImageWrite:
+        case spv::Op::OpImageWrite:
             builtin_name = "textureStore";
             use_level_of_detail_suffix = false;
             if (arg_index < num_args) {
@@ -5511,7 +5588,7 @@
         image_operands_mask = inst.GetSingleWordInOperand(arg_index);
         arg_index++;
     }
-    if (arg_index < num_args && (image_operands_mask & SpvImageOperandsBiasMask)) {
+    if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Bias))) {
         if (is_dref_sample) {
             return Fail() << "WGSL does not support depth-reference sampling with "
                              "level-of-detail bias: "
@@ -5524,10 +5601,10 @@
         }
         builtin_name += "Bias";
         args.Push(MakeOperand(inst, arg_index).expr);
-        image_operands_mask ^= SpvImageOperandsBiasMask;
+        image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Bias);
         arg_index++;
     }
-    if (arg_index < num_args && (image_operands_mask & SpvImageOperandsLodMask)) {
+    if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Lod))) {
         if (use_level_of_detail_suffix) {
             builtin_name += "Level";
         }
@@ -5553,15 +5630,16 @@
             args.Push(lod.expr);
         }
 
-        image_operands_mask ^= SpvImageOperandsLodMask;
+        image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Lod);
         arg_index++;
-    } else if ((opcode == SpvOpImageFetch || opcode == SpvOpImageRead) &&
+    } else if ((op == spv::Op::OpImageFetch || op == spv::Op::OpImageRead) &&
                !texture_type->IsAnyOf<DepthMultisampledTexture, MultisampledTexture>()) {
         // textureLoad requires an explicit level-of-detail parameter for
         // non-multisampled texture types.
         args.Push(parser_impl_.MakeNullValue(ty_.I32()));
     }
-    if (arg_index + 1 < num_args && (image_operands_mask & SpvImageOperandsGradMask)) {
+    if (arg_index + 1 < num_args &&
+        (image_operands_mask & uint32_t(spv::ImageOperandsMask::Grad))) {
         if (is_dref_sample) {
             return Fail() << "WGSL does not support depth-reference sampling with "
                              "explicit gradient: "
@@ -5575,11 +5653,12 @@
         builtin_name += "Grad";
         args.Push(MakeOperand(inst, arg_index).expr);
         args.Push(MakeOperand(inst, arg_index + 1).expr);
-        image_operands_mask ^= SpvImageOperandsGradMask;
+        image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Grad);
         arg_index += 2;
     }
-    if (arg_index < num_args && (image_operands_mask & SpvImageOperandsConstOffsetMask)) {
-        if (!IsImageSamplingOrGatherOrDrefGather(opcode)) {
+    if (arg_index < num_args &&
+        (image_operands_mask & uint32_t(spv::ImageOperandsMask::ConstOffset))) {
+        if (!IsImageSamplingOrGatherOrDrefGather(op)) {
             return Fail() << "ConstOffset is only permitted for sampling, gather, or "
                              "depth-reference gather operations: "
                           << inst.PrettyPrint();
@@ -5596,13 +5675,13 @@
         }
 
         args.Push(ToSignedIfUnsigned(MakeOperand(inst, arg_index)).expr);
-        image_operands_mask ^= SpvImageOperandsConstOffsetMask;
+        image_operands_mask ^= uint32_t(spv::ImageOperandsMask::ConstOffset);
         arg_index++;
     }
-    if (arg_index < num_args && (image_operands_mask & SpvImageOperandsSampleMask)) {
+    if (arg_index < num_args && (image_operands_mask & uint32_t(spv::ImageOperandsMask::Sample))) {
         // TODO(dneto): only permitted with ImageFetch
         args.Push(ToI32(MakeOperand(inst, arg_index)).expr);
-        image_operands_mask ^= SpvImageOperandsSampleMask;
+        image_operands_mask ^= uint32_t(spv::ImageOperandsMask::Sample);
         arg_index++;
     }
     if (image_operands_mask) {
@@ -5642,7 +5721,7 @@
         // Construct a 4-element vector with the result from the builtin in the
         // first component.
         if (texture_type->IsAnyOf<DepthTexture, DepthMultisampledTexture>()) {
-            if (is_non_dref_sample || (opcode == SpvOpImageFetch)) {
+            if (is_non_dref_sample || (op == spv::Op::OpImageFetch)) {
                 value = builder_.Construct(Source{},
                                            result_type->Build(builder_),  // a vec4
                                            utils::Vector{
@@ -5658,7 +5737,7 @@
         // result type. Compare the SPIR-V image's sampled component type with the
         // component of the result type of the SPIR-V instruction.
         auto* spirv_image_type = parser_impl_.GetSpirvTypeForHandleMemoryObjectDeclaration(*image);
-        if (!spirv_image_type || (spirv_image_type->opcode() != SpvOpTypeImage)) {
+        if (!spirv_image_type || (opcode(spirv_image_type) != spv::Op::OpTypeImage)) {
             return Fail() << "invalid image type for image memory object declaration "
                           << image->PrettyPrint();
         }
@@ -5670,7 +5749,7 @@
             value =
                 create<ast::BitcastExpression>(Source{}, result_type->Build(builder_), call_expr);
         }
-        if (!expected_component_type->Is<F32>() && IsSampledImageAccess(opcode)) {
+        if (!expected_component_type->Is<F32>() && IsSampledImageAccess(op)) {
             // WGSL permits sampled image access only on float textures.
             // Reject this case in the SPIR-V reader, at least until SPIR-V validation
             // catches up with this rule and can reject it earlier in the workflow.
@@ -5697,10 +5776,10 @@
         return false;
     }
 
-    const auto opcode = inst.opcode();
-    switch (opcode) {
-        case SpvOpImageQuerySize:
-        case SpvOpImageQuerySizeLod: {
+    const auto op = opcode(inst);
+    switch (op) {
+        case spv::Op::OpImageQuerySize:
+        case spv::Op::OpImageQuerySizeLod: {
             ExpressionList exprs;
             // Invoke textureDimensions.
             // If the texture is arrayed, combine with the result from
@@ -5708,7 +5787,7 @@
             auto* dims_ident = create<ast::IdentifierExpression>(
                 Source{}, builder_.Symbols().Register("textureDimensions"));
             ExpressionList dims_args{GetImageExpression(inst)};
-            if (opcode == SpvOpImageQuerySizeLod) {
+            if (op == spv::Op::OpImageQuerySizeLod) {
                 dims_args.Push(MakeOperand(inst, 1).expr);
             }
             const ast::Expression* dims_call =
@@ -5745,13 +5824,13 @@
 
             return EmitConstDefOrWriteToHoistedVar(inst, expr);
         }
-        case SpvOpImageQueryLod:
+        case spv::Op::OpImageQueryLod:
             return Fail() << "WGSL does not support querying the level of detail of an image: "
                           << inst.PrettyPrint();
-        case SpvOpImageQueryLevels:
-        case SpvOpImageQuerySamples: {
+        case spv::Op::OpImageQueryLevels:
+        case spv::Op::OpImageQuerySamples: {
             const auto* name =
-                (opcode == SpvOpImageQueryLevels) ? "textureNumLevels" : "textureNumSamples";
+                (op == spv::Op::OpImageQueryLevels) ? "textureNumLevels" : "textureNumSamples";
             auto* levels_ident =
                 create<ast::IdentifierExpression>(Source{}, builder_.Symbols().Register(name));
             const ast::Expression* ast_expr = create<ast::CallExpression>(
@@ -5834,46 +5913,46 @@
         return {};
     };
 
-    switch (inst.opcode()) {
-        case SpvOpAtomicLoad:
+    switch (opcode(inst)) {
+        case spv::Op::OpAtomicLoad:
             return emit_atomic(sem::BuiltinType::kAtomicLoad, {oper(/*ptr*/ 0)});
-        case SpvOpAtomicStore:
+        case spv::Op::OpAtomicStore:
             return emit_atomic(sem::BuiltinType::kAtomicStore,
                                {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicExchange:
+        case spv::Op::OpAtomicExchange:
             return emit_atomic(sem::BuiltinType::kAtomicExchange,
                                {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicCompareExchange:
-        case SpvOpAtomicCompareExchangeWeak:
+        case spv::Op::OpAtomicCompareExchange:
+        case spv::Op::OpAtomicCompareExchangeWeak:
             return emit_atomic(sem::BuiltinType::kAtomicCompareExchangeWeak,
                                {oper(/*ptr*/ 0), /*value*/ oper(5), /*comparator*/ oper(4)});
-        case SpvOpAtomicIIncrement:
+        case spv::Op::OpAtomicIIncrement:
             return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), lit(1)});
-        case SpvOpAtomicIDecrement:
+        case spv::Op::OpAtomicIDecrement:
             return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), lit(1)});
-        case SpvOpAtomicIAdd:
+        case spv::Op::OpAtomicIAdd:
             return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicISub:
+        case spv::Op::OpAtomicISub:
             return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicSMin:
+        case spv::Op::OpAtomicSMin:
             return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicUMin:
+        case spv::Op::OpAtomicUMin:
             return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicSMax:
+        case spv::Op::OpAtomicSMax:
             return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicUMax:
+        case spv::Op::OpAtomicUMax:
             return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicAnd:
+        case spv::Op::OpAtomicAnd:
             return emit_atomic(sem::BuiltinType::kAtomicAnd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicOr:
+        case spv::Op::OpAtomicOr:
             return emit_atomic(sem::BuiltinType::kAtomicOr, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicXor:
+        case spv::Op::OpAtomicXor:
             return emit_atomic(sem::BuiltinType::kAtomicXor, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
-        case SpvOpAtomicFlagTestAndSet:
-        case SpvOpAtomicFlagClear:
-        case SpvOpAtomicFMinEXT:
-        case SpvOpAtomicFMaxEXT:
-        case SpvOpAtomicFAddEXT:
+        case spv::Op::OpAtomicFlagTestAndSet:
+        case spv::Op::OpAtomicFlagClear:
+        case spv::Op::OpAtomicFMinEXT:
+        case spv::Op::OpAtomicFMaxEXT:
+        case spv::Op::OpAtomicFAddEXT:
             return Fail() << "unsupported atomic op: " << inst.PrettyPrint();
 
         default:
@@ -5926,11 +6005,11 @@
                << " prompted by " << inst.PrettyPrint();
     }
     bool is_proj = false;
-    switch (inst.opcode()) {
-        case SpvOpImageSampleProjImplicitLod:
-        case SpvOpImageSampleProjExplicitLod:
-        case SpvOpImageSampleProjDrefImplicitLod:
-        case SpvOpImageSampleProjDrefExplicitLod:
+    switch (opcode(inst)) {
+        case spv::Op::OpImageSampleProjImplicitLod:
+        case spv::Op::OpImageSampleProjExplicitLod:
+        case spv::Op::OpImageSampleProjDrefImplicitLod:
+        case spv::Op::OpImageSampleProjDrefExplicitLod:
             is_proj = true;
             break;
         default:
diff --git a/src/tint/reader/spirv/parser_impl.cc b/src/tint/reader/spirv/parser_impl.cc
index 13b4132..f055f4c 100644
--- a/src/tint/reader/spirv/parser_impl.cc
+++ b/src/tint/reader/spirv/parser_impl.cc
@@ -34,6 +34,65 @@
 
 namespace tint::reader::spirv {
 
+namespace three_sided_patch {
+// This machinery is only used while SPIRV-Tools is in transition before it fully
+// uses the C++11 header spirv.hpp11
+
+/// Typedef for pointer to member function while the API call uses
+/// SpvStorageClass for its second argument.
+typedef uint32_t (
+    spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClass)(uint32_t, SpvStorageClass);
+/// Typedef for pointer to member function while the API call uses
+/// spv::StorageClass for its second argument.
+typedef uint32_t (spvtools::opt::analysis::TypeManager::*PointerFinderSpvStorageClassCpp11)(
+    uint32_t,
+    spv::StorageClass);
+
+/// @param type_manager the SPIRV-Tools optimizer's type manager
+/// @param finder a pointer to member function in the type manager that does the
+/// actual lookup
+/// @param pointee_type_id the ID of the pointee type
+/// @param sc the storage class.  SC can be SpvStorageClass or spv::StorageClass
+/// @returns the ID for a SPIR-V pointer to pointee_type_id in storage class sc
+template <typename FinderType, typename SC>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_manager,
+                           FinderType finder,
+                           uint32_t pointee_type_id,
+                           SC sc);
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClass finder,
+                           uint32_t pointee_type_id,
+                           SpvStorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, sc);
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClass finder,
+                           uint32_t pointee_type_id,
+                           spv::StorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, static_cast<SpvStorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClassCpp11 finder,
+                           uint32_t pointee_type_id,
+                           SpvStorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, static_cast<spv::StorageClass>(sc));
+}
+
+template <>
+uint32_t FindPointerToType(spvtools::opt::analysis::TypeManager* type_mgr,
+                           PointerFinderSpvStorageClassCpp11 finder,
+                           uint32_t pointee_type_id,
+                           spv::StorageClass sc) {
+    return (type_mgr->*finder)(pointee_type_id, sc);
+}
+}  // namespace three_sided_patch
+
 namespace {
 
 // Input SPIR-V needs only to conform to Vulkan 1.1 requirements.
@@ -42,6 +101,17 @@
 // will satisfy SPV_ENV_WEBGPU_0 validation.
 const spv_target_env kInputEnv = SPV_ENV_VULKAN_1_1;
 
+/// @param inst a SPIR-V instruction
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction& inst) {
+    return static_cast<spv::Op>(inst.opcode());
+}
+/// @param inst a SPIR-V instruction pointer
+/// @returns Returns the opcode for an instruciton
+inline spv::Op opcode(const spvtools::opt::Instruction* inst) {
+    return static_cast<spv::Op>(inst->opcode());
+}
+
 // A FunctionTraverser is used to compute an ordering of functions in the
 // module such that callees precede callers.
 class FunctionTraverser {
@@ -70,7 +140,7 @@
         visited_.insert(&f);
         for (const auto& bb : f) {
             for (const auto& inst : bb) {
-                if (inst.opcode() != SpvOpFunctionCall) {
+                if (opcode(inst) != spv::Op::OpFunctionCall) {
                     continue;
                 }
                 const auto* callee = id_to_func_[inst.GetSingleWordInOperand(0)];
@@ -89,17 +159,17 @@
 };
 
 // Returns true if the opcode operates as if its operands are signed integral.
-bool AssumesSignedOperands(SpvOp opcode) {
+bool AssumesSignedOperands(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpSNegate:
-        case SpvOpSDiv:
-        case SpvOpSRem:
-        case SpvOpSMod:
-        case SpvOpSLessThan:
-        case SpvOpSLessThanEqual:
-        case SpvOpSGreaterThan:
-        case SpvOpSGreaterThanEqual:
-        case SpvOpConvertSToF:
+        case spv::Op::OpSNegate:
+        case spv::Op::OpSDiv:
+        case spv::Op::OpSRem:
+        case spv::Op::OpSMod:
+        case spv::Op::OpSLessThan:
+        case spv::Op::OpSLessThanEqual:
+        case spv::Op::OpSGreaterThan:
+        case spv::Op::OpSGreaterThanEqual:
+        case spv::Op::OpConvertSToF:
             return true;
         default:
             break;
@@ -126,15 +196,15 @@
 }
 
 // Returns true if the opcode operates as if its operands are unsigned integral.
-bool AssumesUnsignedOperands(SpvOp opcode) {
+bool AssumesUnsignedOperands(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpUDiv:
-        case SpvOpUMod:
-        case SpvOpULessThan:
-        case SpvOpULessThanEqual:
-        case SpvOpUGreaterThan:
-        case SpvOpUGreaterThanEqual:
-        case SpvOpConvertUToF:
+        case spv::Op::OpUDiv:
+        case spv::Op::OpUMod:
+        case spv::Op::OpULessThan:
+        case spv::Op::OpULessThanEqual:
+        case spv::Op::OpUGreaterThan:
+        case spv::Op::OpUGreaterThanEqual:
+        case spv::Op::OpConvertUToF:
             return true;
         default:
             break;
@@ -163,18 +233,18 @@
 // the signedness of the second operand to match the signedness of the
 // first operand, and it's not one of the OpU* or OpS* instructions.
 // (Those are handled via MakeOperand.)
-bool AssumesSecondOperandSignednessMatchesFirstOperand(SpvOp opcode) {
+bool AssumesSecondOperandSignednessMatchesFirstOperand(spv::Op opcode) {
     switch (opcode) {
         // All the OpI* integer binary operations.
-        case SpvOpIAdd:
-        case SpvOpISub:
-        case SpvOpIMul:
-        case SpvOpIEqual:
-        case SpvOpINotEqual:
+        case spv::Op::OpIAdd:
+        case spv::Op::OpISub:
+        case spv::Op::OpIMul:
+        case spv::Op::OpIEqual:
+        case spv::Op::OpINotEqual:
         // All the bitwise integer binary operations.
-        case SpvOpBitwiseAnd:
-        case SpvOpBitwiseOr:
-        case SpvOpBitwiseXor:
+        case spv::Op::OpBitwiseAnd:
+        case spv::Op::OpBitwiseOr:
+        case spv::Op::OpBitwiseXor:
             return true;
         default:
             break;
@@ -184,24 +254,24 @@
 
 // Returns true if the corresponding WGSL operation requires
 // the signedness of the result to match the signedness of the first operand.
-bool AssumesResultSignednessMatchesFirstOperand(SpvOp opcode) {
+bool AssumesResultSignednessMatchesFirstOperand(spv::Op opcode) {
     switch (opcode) {
-        case SpvOpNot:
-        case SpvOpSNegate:
-        case SpvOpBitCount:
-        case SpvOpBitReverse:
-        case SpvOpSDiv:
-        case SpvOpSMod:
-        case SpvOpSRem:
-        case SpvOpIAdd:
-        case SpvOpISub:
-        case SpvOpIMul:
-        case SpvOpBitwiseAnd:
-        case SpvOpBitwiseOr:
-        case SpvOpBitwiseXor:
-        case SpvOpShiftLeftLogical:
-        case SpvOpShiftRightLogical:
-        case SpvOpShiftRightArithmetic:
+        case spv::Op::OpNot:
+        case spv::Op::OpSNegate:
+        case spv::Op::OpBitCount:
+        case spv::Op::OpBitReverse:
+        case spv::Op::OpSDiv:
+        case spv::Op::OpSMod:
+        case spv::Op::OpSRem:
+        case spv::Op::OpIAdd:
+        case spv::Op::OpISub:
+        case spv::Op::OpIMul:
+        case spv::Op::OpBitwiseAnd:
+        case spv::Op::OpBitwiseOr:
+        case spv::Op::OpBitwiseXor:
+        case spv::Op::OpShiftLeftLogical:
+        case spv::Op::OpShiftRightLogical:
+        case spv::Op::OpShiftRightArithmetic:
             return true;
         default:
             break;
@@ -240,12 +310,12 @@
     if (deco.size() < 1) {
         return false;
     }
-    switch (deco[0]) {
-        case SpvDecorationLocation:
-        case SpvDecorationFlat:
-        case SpvDecorationNoPerspective:
-        case SpvDecorationCentroid:
-        case SpvDecorationSample:
+    switch (static_cast<spv::Decoration>(deco[0])) {
+        case spv::Decoration::Location:
+        case spv::Decoration::Flat:
+        case spv::Decoration::NoPerspective:
+        case spv::Decoration::Centroid:
+        case spv::Decoration::Sample:
             return true;
         default:
             break;
@@ -388,16 +458,16 @@
     const auto& decorations = deco_mgr_->GetDecorationsFor(id, true);
     std::unordered_set<uint32_t> visited;
     for (const auto* inst : decorations) {
-        if (inst->opcode() != SpvOpDecorate) {
+        if (opcode(inst) != spv::Op::OpDecorate) {
             continue;
         }
         // Example: OpDecorate %struct_id Block
         // Example: OpDecorate %array_ty ArrayStride 16
         auto decoration_kind = inst->GetSingleWordInOperand(1);
-        switch (decoration_kind) {
+        switch (static_cast<spv::Decoration>(decoration_kind)) {
             // Restrict and RestrictPointer have no effect in graphics APIs.
-            case SpvDecorationRestrict:
-            case SpvDecorationRestrictPointer:
+            case spv::Decoration::Restrict:
+            case spv::Decoration::RestrictPointer:
                 break;
             default:
                 if (visited.emplace(decoration_kind).second) {
@@ -418,15 +488,15 @@
     std::unordered_set<uint32_t> visited;
     for (const auto* inst : decorations) {
         // Example: OpMemberDecorate %struct_id 1 Offset 16
-        if ((inst->opcode() != SpvOpMemberDecorate) ||
+        if ((opcode(inst) != spv::Op::OpMemberDecorate) ||
             (inst->GetSingleWordInOperand(1) != member_index)) {
             continue;
         }
         auto decoration_kind = inst->GetSingleWordInOperand(2);
-        switch (decoration_kind) {
+        switch (static_cast<spv::Decoration>(decoration_kind)) {
             // Restrict and RestrictPointer have no effect in graphics APIs.
-            case SpvDecorationRestrict:
-            case SpvDecorationRestrictPointer:
+            case spv::Decoration::Restrict:
+            case spv::Decoration::RestrictPointer:
                 break;
             default:
                 if (visited.emplace(decoration_kind).second) {
@@ -458,8 +528,8 @@
         Fail() << "malformed SPIR-V decoration: it's empty";
         return {};
     }
-    switch (decoration[0]) {
-        case SpvDecorationOffset:
+    switch (static_cast<spv::Decoration>(decoration[0])) {
+        case spv::Decoration::Offset:
             if (decoration.size() != 2) {
                 Fail() << "malformed Offset decoration: expected 1 literal operand, has "
                        << decoration.size() - 1 << ": member " << member_index << " of "
@@ -469,24 +539,24 @@
             return {
                 builder_.MemberOffset(Source{}, AInt(decoration[1])),
             };
-        case SpvDecorationNonReadable:
+        case spv::Decoration::NonReadable:
             // WGSL doesn't have a member decoration for this.  Silently drop it.
             return {};
-        case SpvDecorationNonWritable:
+        case spv::Decoration::NonWritable:
             // WGSL doesn't have a member decoration for this.
             return {};
-        case SpvDecorationColMajor:
+        case spv::Decoration::ColMajor:
             // WGSL only supports column major matrices.
             return {};
-        case SpvDecorationRelaxedPrecision:
+        case spv::Decoration::RelaxedPrecision:
             // WGSL doesn't support relaxed precision.
             return {};
-        case SpvDecorationRowMajor:
+        case spv::Decoration::RowMajor:
             Fail() << "WGSL does not support row-major matrices: can't "
                       "translate member "
                    << member_index << " of " << ShowType(struct_type_id);
             return {};
-        case SpvDecorationMatrixStride: {
+        case spv::Decoration::MatrixStride: {
             if (decoration.size() != 2) {
                 Fail() << "malformed MatrixStride decoration: expected 1 literal "
                           "operand, has "
@@ -589,15 +659,15 @@
         [this, &in_op_line_scope, &op_line_source,
          &instruction_number](const spvtools::opt::Instruction* inst) {
             ++instruction_number.line;
-            switch (inst->opcode()) {
-                case SpvOpLine:
+            switch (opcode(inst)) {
+                case spv::Op::OpLine:
                     in_op_line_scope = true;
                     // TODO(dneto): This ignores the File ID (operand 0), since the Tint
                     // Source concept doesn't represent that.
                     op_line_source.line = inst->GetSingleWordInOperand(1);
                     op_line_source.column = inst->GetSingleWordInOperand(2);
                     break;
-                case SpvOpNoLine:
+                case spv::Op::OpNoLine:
                     in_op_line_scope = false;
                     break;
                 default:
@@ -670,12 +740,12 @@
 }
 
 bool ParserImpl::IsGlslExtendedInstruction(const spvtools::opt::Instruction& inst) const {
-    return (inst.opcode() == SpvOpExtInst) &&
+    return (opcode(inst) == spv::Op::OpExtInst) &&
            (glsl_std_450_imports_.count(inst.GetSingleWordInOperand(0)) > 0);
 }
 
 bool ParserImpl::IsIgnoredExtendedInstruction(const spvtools::opt::Instruction& inst) const {
-    return (inst.opcode() == SpvOpExtInst) &&
+    return (opcode(inst) == spv::Op::OpExtInst) &&
            (ignored_imports_.count(inst.GetSingleWordInOperand(0)) > 0);
 }
 
@@ -722,15 +792,15 @@
 
     // Register names from OpName and OpMemberName
     for (const auto& inst : module_->debugs2()) {
-        switch (inst.opcode()) {
-            case SpvOpName: {
+        switch (opcode(inst)) {
+            case spv::Op::OpName: {
                 const auto name = inst.GetInOperand(1).AsString();
                 if (!name.empty()) {
                     namer_.SuggestSanitizedName(inst.GetSingleWordInOperand(0), name);
                 }
                 break;
             }
-            case SpvOpMemberName: {
+            case spv::Op::OpMemberName: {
                 const auto name = inst.GetInOperand(2).AsString();
                 if (!name.empty()) {
                     namer_.SuggestSanitizedMemberName(inst.GetSingleWordInOperand(0),
@@ -745,7 +815,7 @@
 
     // Fill in struct member names, and disambiguate them.
     for (const auto* type_inst : module_->GetTypes()) {
-        if (type_inst->opcode() == SpvOpTypeStruct) {
+        if (opcode(type_inst) == spv::Op::OpTypeStruct) {
             namer_.ResolveMemberNamesForStruct(type_inst->result_id(), type_inst->NumInOperands());
         }
     }
@@ -779,13 +849,13 @@
 bool ParserImpl::RegisterWorkgroupSizeBuiltin() {
     WorkgroupSizeInfo& info = workgroup_size_builtin_;
     for (const spvtools::opt::Instruction& inst : module_->annotations()) {
-        if (inst.opcode() != SpvOpDecorate) {
+        if (opcode(inst) != spv::Op::OpDecorate) {
             continue;
         }
-        if (inst.GetSingleWordInOperand(1) != SpvDecorationBuiltIn) {
+        if (inst.GetSingleWordInOperand(1) != uint32_t(spv::Decoration::BuiltIn)) {
             continue;
         }
-        if (inst.GetSingleWordInOperand(2) != SpvBuiltInWorkgroupSize) {
+        if (inst.GetSingleWordInOperand(2) != uint32_t(spv::BuiltIn::WorkgroupSize)) {
             continue;
         }
         info.id = inst.GetSingleWordInOperand(0);
@@ -806,8 +876,8 @@
     // const-expr yet, so avoid supporting OpSpecConstantOp here.
     // TODO(dneto): See https://github.com/gpuweb/gpuweb/issues/1272 for WGSL
     // const_expr proposals.
-    if ((composite_def->opcode() != SpvOpSpecConstantComposite &&
-         composite_def->opcode() != SpvOpConstantComposite)) {
+    if ((opcode(composite_def) != spv::Op::OpSpecConstantComposite &&
+         opcode(composite_def) != spv::Op::OpConstantComposite)) {
         return Fail() << "Invalid WorkgroupSize builtin.  Expected 3-element "
                          "OpSpecConstantComposite or OpConstantComposite:  "
                       << composite_def->PrettyPrint();
@@ -822,7 +892,8 @@
                                            int index) -> bool {
         const auto id = composite_def->GetSingleWordInOperand(static_cast<uint32_t>(index));
         const auto* def = def_use_mgr_->GetDef(id);
-        if (!def || (def->opcode() != SpvOpSpecConstant && def->opcode() != SpvOpConstant) ||
+        if (!def ||
+            (opcode(def) != spv::Op::OpSpecConstant && opcode(def) != spv::Op::OpConstant) ||
             (def->NumInOperands() != 1)) {
             return Fail() << "invalid component " << index << " of workgroupsize "
                           << (def ? def->PrettyPrint() : std::string("no definition"));
@@ -842,8 +913,8 @@
     // decorations.
     std::unordered_map<uint32_t, GridSize> local_size;
     for (const spvtools::opt::Instruction& inst : module_->execution_modes()) {
-        auto mode = static_cast<SpvExecutionMode>(inst.GetSingleWordInOperand(1));
-        if (mode == SpvExecutionModeLocalSize) {
+        auto mode = static_cast<spv::ExecutionMode>(inst.GetSingleWordInOperand(1));
+        if (mode == spv::ExecutionMode::LocalSize) {
             if (inst.NumInOperands() != 5) {
                 // This won't even get past SPIR-V binary parsing.
                 return Fail() << "invalid LocalSize execution mode: " << inst.PrettyPrint();
@@ -856,7 +927,7 @@
     }
 
     for (const spvtools::opt::Instruction& entry_point : module_->entry_points()) {
-        const auto stage = SpvExecutionModel(entry_point.GetSingleWordInOperand(0));
+        const auto stage = spv::ExecutionModel(entry_point.GetSingleWordInOperand(0));
         const uint32_t function_id = entry_point.GetSingleWordInOperand(1);
 
         const std::string ep_name = entry_point.GetOperand(2).AsString();
@@ -886,11 +957,11 @@
         for (unsigned iarg = 3; iarg < entry_point.NumInOperands(); iarg++) {
             const uint32_t var_id = entry_point.GetSingleWordInOperand(iarg);
             if (const auto* var_inst = def_use_mgr_->GetDef(var_id)) {
-                switch (SpvStorageClass(var_inst->GetSingleWordInOperand(0))) {
-                    case SpvStorageClassInput:
+                switch (spv::StorageClass(var_inst->GetSingleWordInOperand(0))) {
+                    case spv::StorageClass::Input:
                         inputs.Add(var_id);
                         break;
-                    case SpvStorageClassOutput:
+                    case spv::StorageClass::Output:
                         outputs.Add(var_id);
                         break;
                     default:
@@ -1036,7 +1107,7 @@
     *array_stride = 0;  // Implicit stride case.
     const auto type_id = type_mgr_->GetId(spv_type);
     for (auto& decoration : this->GetDecorationsFor(type_id)) {
-        if (decoration.size() == 2 && decoration[0] == SpvDecorationArrayStride) {
+        if (decoration.size() == 2 && decoration[0] == uint32_t(spv::Decoration::ArrayStride)) {
             const auto stride = decoration[1];
             if (stride == 0) {
                 return Fail() << "invalid array type ID " << type_id << ": ArrayStride can't be 0";
@@ -1057,9 +1128,9 @@
     auto struct_decorations = this->GetDecorationsFor(type_id);
     if (struct_decorations.size() == 1) {
         const auto decoration = struct_decorations[0][0];
-        if (decoration == SpvDecorationBufferBlock) {
+        if (decoration == uint32_t(spv::Decoration::BufferBlock)) {
             remap_buffer_block_type_.insert(type_id);
-        } else if (decoration != SpvDecorationBlock) {
+        } else if (decoration != uint32_t(spv::Decoration::Block)) {
             Fail() << "struct with ID " << type_id
                    << " has unrecognized decoration: " << int(decoration);
         }
@@ -1097,21 +1168,21 @@
                 Fail() << "malformed SPIR-V decoration: it's empty";
                 return nullptr;
             }
-            if ((decoration[0] == SpvDecorationBuiltIn) && (decoration.size() > 1)) {
-                switch (decoration[1]) {
-                    case SpvBuiltInPosition:
+            if ((decoration[0] == uint32_t(spv::Decoration::BuiltIn)) && (decoration.size() > 1)) {
+                switch (static_cast<spv::BuiltIn>(decoration[1])) {
+                    case spv::BuiltIn::Position:
                         // Record this built-in variable specially.
                         builtin_position_.struct_type_id = type_id;
                         builtin_position_.position_member_index = member_index;
                         builtin_position_.position_member_type_id = member_type_id;
                         create_ast_member = false;  // Not part of the WGSL structure.
                         break;
-                    case SpvBuiltInPointSize:  // not supported in WGSL, but ignore
+                    case spv::BuiltIn::PointSize:  // not supported in WGSL, but ignore
                         builtin_position_.pointsize_member_index = member_index;
                         create_ast_member = false;  // Not part of the WGSL structure.
                         break;
-                    case SpvBuiltInClipDistance:    // not supported in WGSL
-                    case SpvBuiltInCullDistance:    // not supported in WGSL
+                    case spv::BuiltIn::ClipDistance:  // not supported in WGSL
+                    case spv::BuiltIn::CullDistance:  // not supported in WGSL
                         create_ast_member = false;  // Not part of the WGSL structure.
                         break;
                     default:
@@ -1131,7 +1202,7 @@
             if (IsPipelineDecoration(decoration)) {
                 // IO decorations are handled when emitting the entry point.
                 continue;
-            } else if (decoration[0] == SpvDecorationNonWritable) {
+            } else if (decoration[0] == uint32_t(spv::Decoration::NonWritable)) {
                 // WGSL doesn't represent individual members as non-writable. Instead,
                 // apply the ReadOnly access control to the containing struct if all
                 // the members are non-writable.
@@ -1193,12 +1264,12 @@
                                     const spvtools::opt::analysis::Pointer*) {
     const auto* inst = def_use_mgr_->GetDef(type_id);
     const auto pointee_type_id = inst->GetSingleWordInOperand(1);
-    const auto storage_class = SpvStorageClass(inst->GetSingleWordInOperand(0));
+    const auto storage_class = spv::StorageClass(inst->GetSingleWordInOperand(0));
 
     if (pointee_type_id == builtin_position_.struct_type_id) {
         builtin_position_.pointer_type_id = type_id;
         // Pipeline IO builtins map to private variables.
-        builtin_position_.storage_class = SpvStorageClassPrivate;
+        builtin_position_.storage_class = spv::StorageClass::Private;
         return nullptr;
     }
     auto* ast_elem_ty = ConvertType(pointee_type_id, PtrAs::Ptr);
@@ -1244,21 +1315,21 @@
     // in WGSL. In particular, exclude user-defined pipeline IO in a
     // block-decorated struct.
     for (const auto& type_or_value : module_->types_values()) {
-        if (type_or_value.opcode() != SpvOpVariable) {
+        if (opcode(type_or_value) != spv::Op::OpVariable) {
             continue;
         }
         const auto& var = type_or_value;
-        const auto spirv_storage_class = SpvStorageClass(var.GetSingleWordInOperand(0));
-        if ((spirv_storage_class != SpvStorageClassStorageBuffer) &&
-            (spirv_storage_class != SpvStorageClassUniform)) {
+        const auto spirv_storage_class = spv::StorageClass(var.GetSingleWordInOperand(0));
+        if ((spirv_storage_class != spv::StorageClass::StorageBuffer) &&
+            (spirv_storage_class != spv::StorageClass::Uniform)) {
             continue;
         }
         const auto* ptr_type = def_use_mgr_->GetDef(var.type_id());
-        if (ptr_type->opcode() != SpvOpTypePointer) {
+        if (opcode(ptr_type) != spv::Op::OpTypePointer) {
             return Fail() << "OpVariable type expected to be a pointer: " << var.PrettyPrint();
         }
         const auto* store_type = def_use_mgr_->GetDef(ptr_type->GetSingleWordInOperand(1));
-        if (store_type->opcode() == SpvOpTypeStruct) {
+        if (opcode(store_type) == spv::Op::OpTypeStruct) {
             struct_types_for_buffers_.insert(store_type->result_id());
         } else {
             Fail() << "WGSL does not support arrays of buffers: " << var.PrettyPrint();
@@ -1276,7 +1347,8 @@
     // Manufacture a type for the gl_Position variable if we have to.
     if ((builtin_position_.struct_type_id != 0) &&
         (builtin_position_.position_member_pointer_type_id == 0)) {
-        builtin_position_.position_member_pointer_type_id = type_mgr_->FindPointerToType(
+        builtin_position_.position_member_pointer_type_id = three_sided_patch::FindPointerToType(
+            type_mgr_, &spvtools::opt::analysis::TypeManager::FindPointerToType,
             builtin_position_.position_member_type_id, builtin_position_.storage_class);
         ConvertType(builtin_position_.position_member_pointer_type_id);
     }
@@ -1290,13 +1362,13 @@
     for (auto& inst : module_->types_values()) {
         if (const auto* result_type = type_mgr_->GetType(inst.type_id())) {
             if (result_type->AsPointer()) {
-                switch (inst.opcode()) {
-                    case SpvOpVariable:
+                switch (opcode(inst)) {
+                    case spv::Op::OpVariable:
                         // This is the only valid case.
                         break;
-                    case SpvOpUndef:
+                    case spv::Op::OpUndef:
                         return Fail() << "undef pointer is not valid: " << inst.PrettyPrint();
-                    case SpvOpConstantNull:
+                    case spv::Op::OpConstantNull:
                         return Fail() << "null pointer is not valid: " << inst.PrettyPrint();
                     default:
                         return Fail()
@@ -1319,15 +1391,15 @@
         const Type* ast_type = nullptr;
         ast::LiteralExpression* ast_expr = nullptr;
 
-        switch (inst.opcode()) {
-            case SpvOpSpecConstantTrue:
-            case SpvOpSpecConstantFalse: {
+        switch (opcode(inst)) {
+            case spv::Op::OpSpecConstantTrue:
+            case spv::Op::OpSpecConstantFalse: {
                 ast_type = ConvertType(inst.type_id());
                 ast_expr = create<ast::BoolLiteralExpression>(
-                    Source{}, inst.opcode() == SpvOpSpecConstantTrue);
+                    Source{}, opcode(inst) == spv::Op::OpSpecConstantTrue);
                 break;
             }
-            case SpvOpSpecConstant: {
+            case spv::Op::OpSpecConstant: {
                 ast_type = ConvertType(inst.type_id());
                 const uint32_t literal_value = inst.GetSingleWordInOperand(0);
                 ast_expr = Switch(
@@ -1362,7 +1434,7 @@
         if (ast_type && ast_expr) {
             AttributeList spec_id_decos;
             for (const auto& deco : GetDecorationsFor(inst.result_id())) {
-                if ((deco.size() == 2) && (deco[0] == SpvDecorationSpecId)) {
+                if ((deco.size() == 2) && (deco[0] == uint32_t(spv::Decoration::SpecId))) {
                     const uint32_t id = deco[1];
                     if (id > 65535) {
                         return Fail() << "SpecId too large. WGSL override IDs must be "
@@ -1430,16 +1502,16 @@
         return false;
     }
     for (const auto& type_or_value : module_->types_values()) {
-        if (type_or_value.opcode() != SpvOpVariable) {
+        if (opcode(type_or_value) != spv::Op::OpVariable) {
             continue;
         }
         const auto& var = type_or_value;
-        const auto spirv_storage_class = SpvStorageClass(var.GetSingleWordInOperand(0));
+        const auto spirv_storage_class = spv::StorageClass(var.GetSingleWordInOperand(0));
 
         uint32_t type_id = var.type_id();
         if ((type_id == builtin_position_.pointer_type_id) &&
-            ((spirv_storage_class == SpvStorageClassInput) ||
-             (spirv_storage_class == SpvStorageClassOutput))) {
+            ((spirv_storage_class == spv::StorageClass::Input) ||
+             (spirv_storage_class == spv::StorageClass::Output))) {
             // Skip emitting gl_PerVertex.
             builtin_position_.per_vertex_var_id = var.result_id();
             builtin_position_.per_vertex_var_init_id =
@@ -1464,7 +1536,7 @@
             return false;
         }
         const Type* ast_type = nullptr;
-        if (spirv_storage_class == SpvStorageClassUniformConstant) {
+        if (spirv_storage_class == spv::StorageClass::UniformConstant) {
             // These are opaque handles: samplers or textures
             ast_type = GetTypeForHandleVar(var);
             if (!ast_type) {
@@ -1509,9 +1581,9 @@
         if (builtin_position_.per_vertex_var_init_id) {
             // The initializer is complex.
             const auto* init = def_use_mgr_->GetDef(builtin_position_.per_vertex_var_init_id);
-            switch (init->opcode()) {
-                case SpvOpConstantComposite:
-                case SpvOpSpecConstantComposite:
+            switch (opcode(init)) {
+                case spv::Op::OpConstantComposite:
+                case spv::Op::OpSpecConstantComposite:
                     ast_initializer =
                         MakeConstantExpression(
                             init->GetSingleWordInOperand(builtin_position_.position_member_index))
@@ -1541,15 +1613,15 @@
 // @returns the IntConstant for the size of the array, or nullptr
 const spvtools::opt::analysis::IntConstant* ParserImpl::GetArraySize(uint32_t var_id) {
     auto* var = def_use_mgr_->GetDef(var_id);
-    if (!var || var->opcode() != SpvOpVariable) {
+    if (!var || opcode(var) != spv::Op::OpVariable) {
         return nullptr;
     }
     auto* ptr_type = def_use_mgr_->GetDef(var->type_id());
-    if (!ptr_type || ptr_type->opcode() != SpvOpTypePointer) {
+    if (!ptr_type || opcode(ptr_type) != spv::Op::OpTypePointer) {
         return nullptr;
     }
     auto* array_type = def_use_mgr_->GetDef(ptr_type->GetSingleWordInOperand(1));
-    if (!array_type || array_type->opcode() != SpvOpTypeArray) {
+    if (!array_type || opcode(array_type) != spv::Op::OpTypeArray) {
         return nullptr;
     }
     auto* size = constant_mgr_->FindDeclaredConstant(array_type->GetSingleWordInOperand(1));
@@ -1632,23 +1704,23 @@
         if (deco.empty()) {
             return Fail() << "malformed decoration on ID " << id << ": it is empty";
         }
-        if (deco[0] == SpvDecorationBuiltIn) {
+        if (deco[0] == uint32_t(spv::Decoration::BuiltIn)) {
             if (deco.size() == 1) {
                 return Fail() << "malformed BuiltIn decoration on ID " << id << ": has no operand";
             }
-            const auto spv_builtin = static_cast<SpvBuiltIn>(deco[1]);
+            const auto spv_builtin = static_cast<spv::BuiltIn>(deco[1]);
             switch (spv_builtin) {
-                case SpvBuiltInPointSize:
+                case spv::BuiltIn::PointSize:
                     special_builtins_[id] = spv_builtin;
                     return false;  // This is not an error
-                case SpvBuiltInSampleId:
-                case SpvBuiltInVertexIndex:
-                case SpvBuiltInInstanceIndex:
-                case SpvBuiltInLocalInvocationId:
-                case SpvBuiltInLocalInvocationIndex:
-                case SpvBuiltInGlobalInvocationId:
-                case SpvBuiltInWorkgroupId:
-                case SpvBuiltInNumWorkgroups:
+                case spv::BuiltIn::SampleId:
+                case spv::BuiltIn::VertexIndex:
+                case spv::BuiltIn::InstanceIndex:
+                case spv::BuiltIn::LocalInvocationId:
+                case spv::BuiltIn::LocalInvocationIndex:
+                case spv::BuiltIn::GlobalInvocationId:
+                case spv::BuiltIn::WorkgroupId:
+                case spv::BuiltIn::NumWorkgroups:
                     // The SPIR-V variable may signed (because GLSL requires signed for
                     // some of these), but WGSL requires unsigned.  Handle specially
                     // so we always perform the conversion at load and store.
@@ -1660,7 +1732,7 @@
                         }
                     }
                     break;
-                case SpvBuiltInSampleMask: {
+                case spv::BuiltIn::SampleMask: {
                     // In SPIR-V this is used for both input and output variable.
                     // The SPIR-V variable has store type of array of integer scalar,
                     // either signed or unsigned.
@@ -1691,14 +1763,14 @@
         if (transfer_pipeline_io && IsPipelineDecoration(deco)) {
             non_builtin_pipeline_decorations.push_back(deco);
         }
-        if (deco[0] == SpvDecorationDescriptorSet) {
+        if (deco[0] == uint32_t(spv::Decoration::DescriptorSet)) {
             if (deco.size() == 1) {
                 return Fail() << "malformed DescriptorSet decoration on ID " << id
                               << ": has no operand";
             }
             decorations->Push(builder_.Group(Source{}, AInt(deco[1])));
         }
-        if (deco[0] == SpvDecorationBinding) {
+        if (deco[0] == uint32_t(spv::Decoration::Binding)) {
             if (deco.size() == 1) {
                 return Fail() << "malformed Binding decoration on ID " << id << ": has no operand";
             }
@@ -1756,8 +1828,8 @@
 
     for (const auto& deco : decorations) {
         TINT_ASSERT(Reader, deco.size() > 0);
-        switch (deco[0]) {
-            case SpvDecorationLocation:
+        switch (static_cast<spv::Decoration>(deco[0])) {
+            case spv::Decoration::Location:
                 if (deco.size() != 2) {
                     return Fail() << "malformed Location decoration on ID requires one "
                                      "literal operand";
@@ -1768,24 +1840,24 @@
                     type = ast::InterpolationType::kFlat;
                 }
                 break;
-            case SpvDecorationFlat:
+            case spv::Decoration::Flat:
                 type = ast::InterpolationType::kFlat;
                 break;
-            case SpvDecorationNoPerspective:
+            case spv::Decoration::NoPerspective:
                 if (store_type->IsIntegerScalarOrVector()) {
                     // This doesn't capture the array or struct case.
                     return Fail() << "NoPerspective is invalid on integral IO";
                 }
                 type = ast::InterpolationType::kLinear;
                 break;
-            case SpvDecorationCentroid:
+            case spv::Decoration::Centroid:
                 if (store_type->IsIntegerScalarOrVector()) {
                     // This doesn't capture the array or struct case.
                     return Fail() << "Centroid interpolation sampling is invalid on integral IO";
                 }
                 sampling = ast::InterpolationSampling::kCentroid;
                 break;
-            case SpvDecorationSample:
+            case spv::Decoration::Sample:
                 if (store_type->IsIntegerScalarOrVector()) {
                     // This doesn't capture the array or struct case.
                     return Fail() << "Sample interpolation sampling is invalid on integral IO";
@@ -1801,7 +1873,7 @@
         !ast::HasAttribute<ast::LocationAttribute>(*attributes)) {
         // WGSL requires that '@interpolate(flat)' needs to be paired with '@location', however
         // SPIR-V requires all fragment shader integer Inputs are 'flat'. If the decorations do not
-        // contain a SpvDecorationLocation, then make this perspective.
+        // contain a spv::Decoration::Location, then make this perspective.
         type = ast::InterpolationType::kPerspective;
     }
 
@@ -1825,7 +1897,7 @@
     if (!inst) {
         return false;
     }
-    if (inst->opcode() == SpvOpUndef) {
+    if (opcode(inst) == spv::Op::OpUndef) {
         return true;
     }
     return nullptr != constant_mgr_->FindDeclaredConstant(id);
@@ -1880,13 +1952,13 @@
         return {};
     }
 
-    switch (inst->opcode()) {
-        case SpvOpUndef:  // Remap undef to null.
-        case SpvOpConstantNull:
+    switch (opcode(inst)) {
+        case spv::Op::OpUndef:  // Remap undef to null.
+        case spv::Op::OpConstantNull:
             return {original_ast_type, MakeNullValue(original_ast_type)};
-        case SpvOpConstantTrue:
-        case SpvOpConstantFalse:
-        case SpvOpConstant: {
+        case spv::Op::OpConstantTrue:
+        case spv::Op::OpConstantFalse:
+        case spv::Op::OpConstant: {
             const auto* spirv_const = constant_mgr_->FindDeclaredConstant(id);
             if (spirv_const == nullptr) {
                 Fail() << "ID " << id << " is not a constant";
@@ -1895,7 +1967,7 @@
             return MakeConstantExpressionForScalarSpirvConstant(source, original_ast_type,
                                                                 spirv_const);
         }
-        case SpvOpConstantComposite: {
+        case spv::Op::OpConstantComposite: {
             // Handle vector, matrix, array, and struct
 
             // Generate a composite from explicit components.
@@ -2047,9 +2119,9 @@
         requires_signed = AssumesSignedOperands(extended_opcode);
         requires_unsigned = AssumesUnsignedOperands(extended_opcode);
     } else {
-        const auto opcode = inst.opcode();
-        requires_signed = AssumesSignedOperands(opcode);
-        requires_unsigned = AssumesUnsignedOperands(opcode);
+        const auto op = opcode(inst);
+        requires_signed = AssumesSignedOperands(op);
+        requires_unsigned = AssumesUnsignedOperands(op);
     }
     if (!requires_signed && !requires_unsigned) {
         // No conversion is required, assuming our tables are complete.
@@ -2087,7 +2159,7 @@
                                                            TypedExpression&& second_operand_expr) {
     const Type* target_type = first_operand_type->UnwrapRef();
     if ((target_type != second_operand_expr.type->UnwrapRef()) &&
-        AssumesSecondOperandSignednessMatchesFirstOperand(inst.opcode())) {
+        AssumesSecondOperandSignednessMatchesFirstOperand(opcode(inst))) {
         // Conversion is required.
         return {target_type, create<ast::BitcastExpression>(Source{}, target_type->Build(builder_),
                                                             second_operand_expr.expr)};
@@ -2099,8 +2171,8 @@
 const Type* ParserImpl::ForcedResultType(const spvtools::opt::Instruction& inst,
                                          const Type* first_operand_type) {
     first_operand_type = first_operand_type->UnwrapRef();
-    const auto opcode = inst.opcode();
-    if (AssumesResultSignednessMatchesFirstOperand(opcode)) {
+    const auto op = opcode(inst);
+    if (AssumesResultSignednessMatchesFirstOperand(op)) {
         return first_operand_type;
     }
     if (IsGlslExtendedInstruction(inst)) {
@@ -2239,36 +2311,36 @@
         if (inst == nullptr) {
             return local_fail();
         }
-        switch (inst->opcode()) {
-            case SpvOpFunctionParameter:
-            case SpvOpVariable:
+        switch (opcode(inst)) {
+            case spv::Op::OpFunctionParameter:
+            case spv::Op::OpVariable:
                 // We found the memory object declaration.
                 // Remember it as the answer for the whole path.
                 for (auto iter : visited) {
                     memo_table[iter] = inst;
                 }
                 return inst;
-            case SpvOpLoad:
+            case spv::Op::OpLoad:
                 // Follow the pointer being loaded
                 id = inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpCopyObject:
+            case spv::Op::OpCopyObject:
                 // Follow the object being copied.
                 id = inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpAccessChain:
-            case SpvOpInBoundsAccessChain:
-            case SpvOpPtrAccessChain:
-            case SpvOpInBoundsPtrAccessChain:
+            case spv::Op::OpAccessChain:
+            case spv::Op::OpInBoundsAccessChain:
+            case spv::Op::OpPtrAccessChain:
+            case spv::Op::OpInBoundsPtrAccessChain:
                 // Follow the base pointer.
                 id = inst->GetSingleWordInOperand(0);
                 break;
-            case SpvOpSampledImage:
+            case spv::Op::OpSampledImage:
                 // Follow the image or the sampler, depending on the follow_image
                 // parameter.
                 id = inst->GetSingleWordInOperand(follow_image ? 0 : 1);
                 break;
-            case SpvOpImage:
+            case spv::Op::OpImage:
                 // Follow the sampled image
                 id = inst->GetSingleWordInOperand(0);
                 break;
@@ -2300,7 +2372,7 @@
 
     // Get the SPIR-V handle type.
     const auto* ptr_type = def_use_mgr_->GetDef(var.type_id());
-    if (!ptr_type || (ptr_type->opcode() != SpvOpTypePointer)) {
+    if (!ptr_type || (opcode(ptr_type) != spv::Op::OpTypePointer)) {
         Fail() << "Invalid type for variable or function parameter " << var.PrettyPrint();
         return nullptr;
     }
@@ -2309,18 +2381,18 @@
         Fail() << "Invalid pointer type for variable or function parameter " << var.PrettyPrint();
         return nullptr;
     }
-    switch (raw_handle_type->opcode()) {
-        case SpvOpTypeSampler:
-        case SpvOpTypeImage:
+    switch (opcode(raw_handle_type)) {
+        case spv::Op::OpTypeSampler:
+        case spv::Op::OpTypeImage:
             // The expected cases.
             break;
-        case SpvOpTypeArray:
-        case SpvOpTypeRuntimeArray:
+        case spv::Op::OpTypeArray:
+        case spv::Op::OpTypeRuntimeArray:
             Fail() << "arrays of textures or samplers are not supported in WGSL; can't "
                       "translate variable or function parameter: "
                    << var.PrettyPrint();
             return nullptr;
-        case SpvOpTypeSampledImage:
+        case spv::Op::OpTypeSampledImage:
             Fail() << "WGSL does not support combined image-samplers: " << var.PrettyPrint();
             return nullptr;
         default:
@@ -2358,7 +2430,7 @@
         // In SPIR-V you could statically reference a texture or sampler without
         // using it in a way that gives us a clue on how to declare it.  Look inside
         // the store type to infer a usage.
-        if (raw_handle_type->opcode() == SpvOpTypeSampler) {
+        if (opcode(raw_handle_type) == spv::Op::OpTypeSampler) {
             usage.AddSampler();
         } else {
             // It's a texture.
@@ -2370,7 +2442,7 @@
             const auto sampled_param = raw_handle_type->GetSingleWordInOperand(5);
             const auto format_param = raw_handle_type->GetSingleWordInOperand(6);
             // Only storage images have a format.
-            if ((format_param != SpvImageFormatUnknown) ||
+            if ((format_param != uint32_t(spv::ImageFormat::Unknown)) ||
                 sampled_param == 2 /* without sampler */) {
                 // Get NonWritable and NonReadable attributes of the variable.
                 bool is_nonwritable = false;
@@ -2379,10 +2451,10 @@
                     if (deco.size() != 1) {
                         continue;
                     }
-                    if (deco[0] == SpvDecorationNonWritable) {
+                    if (deco[0] == uint32_t(spv::Decoration::NonWritable)) {
                         is_nonwritable = true;
                     }
-                    if (deco[0] == SpvDecorationNonReadable) {
+                    if (deco[0] == uint32_t(spv::Decoration::NonReadable)) {
                         is_nonreadable = true;
                     }
                 }
@@ -2429,9 +2501,9 @@
         if (image_type->is_arrayed()) {
             // Give a nicer error message here, where we have the offending variable
             // in hand, rather than inside the enum converter.
-            switch (image_type->dim()) {
-                case SpvDim2D:
-                case SpvDimCube:
+            switch (static_cast<spv::Dim>(image_type->dim())) {
+                case spv::Dim::Dim2D:
+                case spv::Dim::Cube:
                     break;
                 default:
                     Fail() << "WGSL arrayed textures must be 2d_array or cube_array: "
@@ -2450,7 +2522,7 @@
         // WGSL textures are always formatted.  Unformatted textures are always
         // sampled.
         if (usage.IsSampledTexture() || usage.IsStorageReadTexture() ||
-            (image_type->format() == SpvImageFormatUnknown)) {
+            (uint32_t(image_type->format()) == uint32_t(spv::ImageFormat::Unknown))) {
             // Make a sampled texture type.
             auto* ast_sampled_component_type =
                 ConvertType(raw_handle_type->GetSingleWordInOperand(0));
@@ -2607,34 +2679,34 @@
     for (const auto* f : topologically_ordered_functions_) {
         for (const auto& bb : *f) {
             for (const auto& inst : bb) {
-                switch (inst.opcode()) {
+                switch (opcode(inst)) {
                         // Single texel reads and writes
 
-                    case SpvOpImageRead:
+                    case spv::Op::OpImageRead:
                         handle_usage_[get_image(inst)].AddStorageReadTexture();
                         break;
-                    case SpvOpImageWrite:
+                    case spv::Op::OpImageWrite:
                         handle_usage_[get_image(inst)].AddStorageWriteTexture();
                         break;
-                    case SpvOpImageFetch:
+                    case spv::Op::OpImageFetch:
                         handle_usage_[get_image(inst)].AddSampledTexture();
                         break;
 
                         // Sampling and gathering from a sampled image.
 
-                    case SpvOpImageSampleImplicitLod:
-                    case SpvOpImageSampleExplicitLod:
-                    case SpvOpImageSampleProjImplicitLod:
-                    case SpvOpImageSampleProjExplicitLod:
-                    case SpvOpImageGather:
+                    case spv::Op::OpImageSampleImplicitLod:
+                    case spv::Op::OpImageSampleExplicitLod:
+                    case spv::Op::OpImageSampleProjImplicitLod:
+                    case spv::Op::OpImageSampleProjExplicitLod:
+                    case spv::Op::OpImageGather:
                         handle_usage_[get_image(inst)].AddSampledTexture();
                         handle_usage_[get_sampler(inst)].AddSampler();
                         break;
-                    case SpvOpImageSampleDrefImplicitLod:
-                    case SpvOpImageSampleDrefExplicitLod:
-                    case SpvOpImageSampleProjDrefImplicitLod:
-                    case SpvOpImageSampleProjDrefExplicitLod:
-                    case SpvOpImageDrefGather:
+                    case spv::Op::OpImageSampleDrefImplicitLod:
+                    case spv::Op::OpImageSampleDrefExplicitLod:
+                    case spv::Op::OpImageSampleProjDrefImplicitLod:
+                    case spv::Op::OpImageSampleProjDrefExplicitLod:
+                    case spv::Op::OpImageDrefGather:
                         // Depth reference access implies usage as a depth texture, which
                         // in turn is a sampled texture.
                         handle_usage_[get_image(inst)].AddDepthTexture();
@@ -2643,29 +2715,29 @@
 
                         // Image queries
 
-                    case SpvOpImageQuerySizeLod:
+                    case spv::Op::OpImageQuerySizeLod:
                         // Vulkan requires Sampled=1 for this. SPIR-V already requires MS=0.
                         handle_usage_[get_image(inst)].AddSampledTexture();
                         break;
-                    case SpvOpImageQuerySize:
+                    case spv::Op::OpImageQuerySize:
                         // Applies to either MS=1 or Sampled=0 or 2.
                         // So we can't force it to be multisampled, or storage image.
                         break;
-                    case SpvOpImageQueryLod:
+                    case spv::Op::OpImageQueryLod:
                         handle_usage_[get_image(inst)].AddSampledTexture();
                         handle_usage_[get_sampler(inst)].AddSampler();
                         break;
-                    case SpvOpImageQueryLevels:
+                    case spv::Op::OpImageQueryLevels:
                         // We can't tell anything more than that it's an image.
                         handle_usage_[get_image(inst)].AddTexture();
                         break;
-                    case SpvOpImageQuerySamples:
+                    case spv::Op::OpImageQuerySamples:
                         handle_usage_[get_image(inst)].AddMultisampledTexture();
                         break;
 
                         // Function calls
 
-                    case SpvOpFunctionCall: {
+                    case spv::Op::OpFunctionCall: {
                         // Propagate handle usages from callee function formal parameters to
                         // the matching caller parameters.  This is where we rely on the
                         // fact that callees have been processed earlier in the flow.
diff --git a/src/tint/reader/spirv/parser_impl.h b/src/tint/reader/spirv/parser_impl.h
index 9965aa4..6e79ecc 100644
--- a/src/tint/reader/spirv/parser_impl.h
+++ b/src/tint/reader/spirv/parser_impl.h
@@ -66,8 +66,8 @@
 
 /// The binary representation of a SPIR-V decoration enum followed by its
 /// operands, if any.
-/// Example:   { SpvDecorationBlock }
-/// Example:   { SpvDecorationArrayStride, 16 }
+/// Example:   { spv::Decoration::Block }
+/// Example:   { spv::Decoration::ArrayStride, 16 }
 using Decoration = std::vector<uint32_t>;
 
 /// DecorationList is a list of decorations
@@ -338,13 +338,13 @@
     /// Returns true when the given instruction is an extended instruction
     /// for GLSL.std.450.
     /// @param inst a SPIR-V instruction
-    /// @returns true if its an SpvOpExtInst for GLSL.std.450
+    /// @returns true if its an spv::Op::ExtInst for GLSL.std.450
     bool IsGlslExtendedInstruction(const spvtools::opt::Instruction& inst) const;
 
     /// Returns true when the given instruction is an extended instruction
     /// from an ignored extended instruction set.
     /// @param inst a SPIR-V instruction
-    /// @returns true if its an SpvOpExtInst for an ignored extended instruction
+    /// @returns true if its an spv::Op::ExtInst for an ignored extended instruction
     bool IsIgnoredExtendedInstruction(const spvtools::opt::Instruction& inst) const;
 
     /// Registers user names for SPIR-V objects, from OpName, and OpMemberName.
@@ -585,7 +585,7 @@
         /// class class.
         uint32_t pointer_type_id = 0;
         /// The SPIR-V address space.
-        SpvStorageClass storage_class = SpvStorageClassOutput;
+        spv::StorageClass storage_class = spv::StorageClass::Output;
         /// The ID of the type of a pointer to the Position member.
         uint32_t position_member_pointer_type_id = 0;
         /// The ID of the gl_PerVertex variable, if it was declared.
@@ -690,7 +690,7 @@
     const spvtools::opt::Instruction* GetInstructionForTest(uint32_t id) const;
 
     /// A map of SPIR-V identifiers to builtins
-    using BuiltInsMap = std::unordered_map<uint32_t, SpvBuiltIn>;
+    using BuiltInsMap = std::unordered_map<uint32_t, spv::BuiltIn>;
 
     /// @returns a map of builtins that should be handled specially by code
     /// generation. Either the builtin does not exist in WGSL, or a type
@@ -699,7 +699,7 @@
 
     /// @param builtin the SPIR-V builtin variable kind
     /// @returns the SPIR-V ID for the variable defining the given builtin, or 0
-    uint32_t IdForSpecialBuiltIn(SpvBuiltIn builtin) const {
+    uint32_t IdForSpecialBuiltIn(spv::BuiltIn builtin) const {
         // Do a linear search.
         for (const auto& entry : special_builtins_) {
             if (entry.second == builtin) {
diff --git a/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc b/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
index 767b333..402cc41 100644
--- a/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
+++ b/src/tint/reader/spirv/parser_impl_convert_member_decoration_test.cc
@@ -31,7 +31,7 @@
 TEST_F(SpvParserTest, ConvertMemberDecoration_OffsetWithoutOperand) {
     auto p = parser(std::vector<uint32_t>{});
 
-    auto result = p->ConvertMemberDecoration(12, 13, nullptr, {SpvDecorationOffset});
+    auto result = p->ConvertMemberDecoration(12, 13, nullptr, {uint32_t(spv::Decoration::Offset)});
     EXPECT_TRUE(result.IsEmpty());
     EXPECT_THAT(p->error(), Eq("malformed Offset decoration: expected 1 literal "
                                "operand, has 0: member 13 of SPIR-V type 12"));
@@ -40,7 +40,8 @@
 TEST_F(SpvParserTest, ConvertMemberDecoration_OffsetWithTooManyOperands) {
     auto p = parser(std::vector<uint32_t>{});
 
-    auto result = p->ConvertMemberDecoration(12, 13, nullptr, {SpvDecorationOffset, 3, 4});
+    auto result =
+        p->ConvertMemberDecoration(12, 13, nullptr, {uint32_t(spv::Decoration::Offset), 3, 4});
     EXPECT_TRUE(result.IsEmpty());
     EXPECT_THAT(p->error(), Eq("malformed Offset decoration: expected 1 literal "
                                "operand, has 2: member 13 of SPIR-V type 12"));
@@ -49,7 +50,7 @@
 TEST_F(SpvParserTest, ConvertMemberDecoration_Offset) {
     auto p = parser(std::vector<uint32_t>{});
 
-    auto result = p->ConvertMemberDecoration(1, 1, nullptr, {SpvDecorationOffset, 8});
+    auto result = p->ConvertMemberDecoration(1, 1, nullptr, {uint32_t(spv::Decoration::Offset), 8});
     ASSERT_FALSE(result.IsEmpty());
     EXPECT_TRUE(result[0]->Is<ast::StructMemberOffsetAttribute>());
     auto* offset_deco = result[0]->As<ast::StructMemberOffsetAttribute>();
@@ -64,7 +65,8 @@
 
     spirv::F32 f32;
     spirv::Matrix matrix(&f32, 2, 2);
-    auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 8});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 8});
     EXPECT_TRUE(result.IsEmpty());
     EXPECT_TRUE(p->error().empty());
 }
@@ -74,7 +76,8 @@
 
     spirv::F32 f32;
     spirv::Matrix matrix(&f32, 2, 2);
-    auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 16});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 16});
     ASSERT_FALSE(result.IsEmpty());
     EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
     auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -88,7 +91,8 @@
 
     spirv::F32 f32;
     spirv::Matrix matrix(&f32, 2, 4);
-    auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 16});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 16});
     EXPECT_TRUE(result.IsEmpty());
     EXPECT_TRUE(p->error().empty());
 }
@@ -98,7 +102,8 @@
 
     spirv::F32 f32;
     spirv::Matrix matrix(&f32, 2, 4);
-    auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 64});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 64});
     ASSERT_FALSE(result.IsEmpty());
     EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
     auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -112,7 +117,8 @@
 
     spirv::F32 f32;
     spirv::Matrix matrix(&f32, 2, 3);
-    auto result = p->ConvertMemberDecoration(1, 1, &matrix, {SpvDecorationMatrixStride, 32});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, &matrix, {uint32_t(spv::Decoration::MatrixStride), 32});
     ASSERT_FALSE(result.IsEmpty());
     EXPECT_TRUE(result[0]->Is<ast::StrideAttribute>());
     auto* stride_deco = result[0]->As<ast::StrideAttribute>();
@@ -127,7 +133,8 @@
     // relaxed precision f32.
     auto p = parser(std::vector<uint32_t>{});
 
-    auto result = p->ConvertMemberDecoration(1, 1, nullptr, {SpvDecorationRelaxedPrecision});
+    auto result =
+        p->ConvertMemberDecoration(1, 1, nullptr, {uint32_t(spv::Decoration::RelaxedPrecision)});
     EXPECT_TRUE(result.IsEmpty());
     EXPECT_TRUE(p->error().empty());
 }
diff --git a/src/tint/reader/spirv/parser_impl_get_decorations_test.cc b/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
index 929ab1f..cbc384f 100644
--- a/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
+++ b/src/tint/reader/spirv/parser_impl_get_decorations_test.cc
@@ -52,7 +52,7 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule());
     auto decorations = p->GetDecorationsFor(10);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationBlock}));
+    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::Block)}));
     EXPECT_TRUE(p->error().empty());
     p->SkipDumpingPending(kSkipReason);
 }
@@ -66,7 +66,7 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule());
     auto decorations = p->GetDecorationsFor(10);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationBlock}));
+    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::Block)}));
     EXPECT_TRUE(p->error().empty());
     p->SkipDumpingPending(kSkipReason);
 }
@@ -80,8 +80,9 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule());
     auto decorations = p->GetDecorationsFor(5);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision},
-                                                  Decoration{SpvDecorationLocation, 7}));
+    EXPECT_THAT(decorations,
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)},
+                                     Decoration{uint32_t(spv::Decoration::Location), 7}));
     EXPECT_TRUE(p->error().empty());
     p->SkipDumpingPending(kSkipReason);
 }
@@ -124,7 +125,8 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
     auto decorations = p->GetDecorationsForMember(10, 0);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+    EXPECT_THAT(decorations,
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
     EXPECT_TRUE(p->error().empty());
     p->SkipDumpingPending(kSkipReason);
 }
@@ -138,7 +140,8 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
     auto decorations = p->GetDecorationsForMember(10, 0);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+    EXPECT_THAT(decorations,
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
     EXPECT_TRUE(p->error().empty());
     p->SkipDumpingPending(kSkipReason);
 }
@@ -154,7 +157,8 @@
   )"));
     EXPECT_TRUE(p->BuildAndParseInternalModule()) << p->error();
     auto decorations = p->GetDecorationsForMember(10, 1);
-    EXPECT_THAT(decorations, UnorderedElementsAre(Decoration{SpvDecorationArrayStride, 12}));
+    EXPECT_THAT(decorations,
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::ArrayStride), 12}));
     EXPECT_TRUE(p->error().empty());
 }
 
@@ -179,11 +183,11 @@
 
     EXPECT_TRUE(p->GetDecorationsForMember(50, 0).empty());
     EXPECT_THAT(p->GetDecorationsForMember(50, 1),
-                UnorderedElementsAre(Decoration{SpvDecorationRelaxedPrecision}));
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::RelaxedPrecision)}));
     EXPECT_THAT(p->GetDecorationsForMember(50, 2),
-                UnorderedElementsAre(Decoration{SpvDecorationColMajor},
-                                     Decoration{SpvDecorationMatrixStride, 8},
-                                     Decoration{SpvDecorationArrayStride, 16}));
+                UnorderedElementsAre(Decoration{uint32_t(spv::Decoration::ColMajor)},
+                                     Decoration{uint32_t(spv::Decoration::MatrixStride), 8},
+                                     Decoration{uint32_t(spv::Decoration::ArrayStride), 16}));
     EXPECT_TRUE(p->error().empty());
 }
 
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 976706a..398c123 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -615,12 +615,11 @@
     Source source;
     if (match(Token::Type::kConst)) {
         use = "'const' declaration";
-    } else if (match(Token::Type::kLet, &source)) {
-        use = "'let' declaration";
-        deprecated(source, "module-scope 'let' has been replaced with 'const'");
     } else if (match(Token::Type::kOverride)) {
         use = "'override' declaration";
         is_overridable = true;
+    } else if (match(Token::Type::kLet, &source)) {
+        return add_error(source, "module-scope 'let' is invalid, use 'const'");
     } else {
         return Failure::kNoMatch;
     }
@@ -1613,29 +1612,6 @@
                                   std::move(attrs.value));                  // attributes
 }
 
-// pipeline_stage
-//   : VERTEX
-//   | FRAGMENT
-//   | COMPUTE
-//
-// TODO(crbug.com/tint/1503): Remove when deprecation period is over.
-Expect<ast::PipelineStage> ParserImpl::expect_pipeline_stage() {
-    auto& t = peek();
-    if (t == "vertex") {
-        next();  // Consume the peek
-        return {ast::PipelineStage::kVertex, t.source()};
-    }
-    if (t == "fragment") {
-        next();  // Consume the peek
-        return {ast::PipelineStage::kFragment, t.source()};
-    }
-    if (t == "compute") {
-        next();  // Consume the peek
-        return {ast::PipelineStage::kCompute, t.source()};
-    }
-    return add_error(peek(), "invalid value for stage attribute");
-}
-
 // interpolation_sample_name
 //   : 'center'
 //   | 'centroid'
@@ -3637,34 +3613,6 @@
         });
     }
 
-    // TODO(crbug.com/tint/1503): Remove when deprecation period is over.
-    if (t == "stage") {
-        return expect_paren_block("stage attribute", [&]() -> Result {
-            auto stage = expect_pipeline_stage();
-            if (stage.errored) {
-                return Failure::kErrored;
-            }
-
-            std::string warning = "remove stage and use @";
-            switch (stage.value) {
-                case ast::PipelineStage::kVertex:
-                    warning += "vertex";
-                    break;
-                case ast::PipelineStage::kFragment:
-                    warning += "fragment";
-                    break;
-                case ast::PipelineStage::kCompute:
-                    warning += "compute";
-                    break;
-                case ast::PipelineStage::kNone:
-                    break;
-            }
-            deprecated(t.source(), warning);
-
-            return create<ast::StageAttribute>(t.source(), stage.value);
-        });
-    }
-
     if (t == "vertex") {
         return create<ast::StageAttribute>(t.source(), ast::PipelineStage::kVertex);
     }
@@ -3826,7 +3774,7 @@
         next();
 
         if (is_reserved(t)) {
-            deprecated(t.source(), "'" + t.to_str() + "' is a reserved keyword");
+            return add_error(t.source(), "'" + t.to_str() + "' is a reserved keyword");
         }
 
         return {t.to_str(), t.source()};
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index fbfe6f6..6c4f14a 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -351,48 +351,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1503): Remove this when @stage is removed
-TEST_F(ParserImplErrorTest, FunctionDeclStageMissingLParen) {
-    EXPECT("@stage vertex) fn f() {}",
-           R"(test.wgsl:1:8 error: expected '(' for stage attribute
-@stage vertex) fn f() {}
-       ^^^^^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageMissingRParen) {
-    EXPECT(
-        "@stage(vertex fn f() {}",
-        R"(test.wgsl:1:2 warning: use of deprecated language feature: remove stage and use @vertex
-@stage(vertex fn f() {}
- ^^^^^
-
-test.wgsl:1:15 error: expected ')' for stage attribute
-@stage(vertex fn f() {}
-              ^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageInvalid) {
-    EXPECT("@stage(x) fn f() {}",
-           R"(test.wgsl:1:8 error: invalid value for stage attribute
-@stage(x) fn f() {}
-       ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, FunctionDeclStageTypeInvalid) {
-    EXPECT("@shader(vertex) fn main() {}",
-           R"(test.wgsl:1:2 error: expected attribute
-@shader(vertex) fn main() {}
- ^^^^^^
-
-test.wgsl:1:8 error: unexpected token
-@shader(vertex) fn main() {}
-       ^
-)");
-}
-
 TEST_F(ParserImplErrorTest, FunctionDeclWorkgroupSizeXInvalid) {
     EXPECT("@workgroup_size() fn f() {}",
            R"(test.wgsl:1:17 error: expected workgroup_size x parameter
@@ -595,94 +553,11 @@
 )");
 }
 
-TEST_F(ParserImplErrorTest, GlobalDeclLetInvalidIdentifier) {
-    EXPECT(
-        "let ^ : i32 = 1;",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let ^ : i32 = 1;
+TEST_F(ParserImplErrorTest, GlobalDeclLet) {
+    EXPECT("let a : i32 = 1;",
+           R"(test.wgsl:1:1 error: module-scope 'let' is invalid, use 'const'
+let a : i32 = 1;
 ^^^
-
-test.wgsl:1:5 error: expected identifier for 'let' declaration
-let ^ : i32 = 1;
-    ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingSemicolon) {
-    EXPECT(
-        "let i : i32 = 1",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : i32 = 1
-^^^
-
-test.wgsl:1:16 error: expected ';' for 'const' declaration
-let i : i32 = 1
-               ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingLParen) {
-    EXPECT(
-        "let i : vec2<i32> = vec2<i32>;",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>;
-^^^
-
-test.wgsl:1:30 error: expected '(' for type initializer
-let i : vec2<i32> = vec2<i32>;
-                             ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetMissingRParen) {
-    EXPECT(
-        "let i : vec2<i32> = vec2<i32>(1., 2.;",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(1., 2.;
-^^^
-
-test.wgsl:1:37 error: expected ')' for type initializer
-let i : vec2<i32> = vec2<i32>(1., 2.;
-                                    ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetBadConstLiteral) {
-    EXPECT(
-        "let i : vec2<i32> = vec2<i32>(!);",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(!);
-^^^
-
-test.wgsl:1:32 error: unable to parse right side of ! expression
-let i : vec2<i32> = vec2<i32>(!);
-                               ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetExprMissingLParen) {
-    EXPECT(
-        "let i : vec2<i32> = vec2<i32> 1, 2);",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32> 1, 2);
-^^^
-
-test.wgsl:1:31 error: expected '(' for type initializer
-let i : vec2<i32> = vec2<i32> 1, 2);
-                              ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclLetExprMissingRParen) {
-    EXPECT(
-        "let i : vec2<i32> = vec2<i32>(1, 2;",
-        R"(test.wgsl:1:1 warning: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-let i : vec2<i32> = vec2<i32>(1, 2;
-^^^
-
-test.wgsl:1:35 error: expected ')' for type initializer
-let i : vec2<i32> = vec2<i32>(1, 2;
-                                  ^
 )");
 }
 
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
index 28b84a9..785d23c 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
@@ -18,33 +18,6 @@
 namespace tint::reader::wgsl {
 namespace {
 
-// TODO(crbug.com/tint/1503): Remove this when @stage is removed
-TEST_F(ParserImplTest, AttributeList_Parses_Stage) {
-    auto p = parser("@workgroup_size(2) @stage(compute)");
-    auto attrs = p->attribute_list();
-    EXPECT_FALSE(p->has_error()) << p->error();
-    EXPECT_FALSE(attrs.errored);
-    EXPECT_TRUE(attrs.matched);
-    ASSERT_EQ(attrs.value.Length(), 2u);
-
-    auto* attr_0 = attrs.value[0]->As<ast::Attribute>();
-    auto* attr_1 = attrs.value[1]->As<ast::Attribute>();
-    ASSERT_NE(attr_0, nullptr);
-    ASSERT_NE(attr_1, nullptr);
-
-    ASSERT_TRUE(attr_0->Is<ast::WorkgroupAttribute>());
-    const ast::Expression* x = attr_0->As<ast::WorkgroupAttribute>()->x;
-    ASSERT_NE(x, nullptr);
-    auto* x_literal = x->As<ast::LiteralExpression>();
-    ASSERT_NE(x_literal, nullptr);
-    ASSERT_TRUE(x_literal->Is<ast::IntLiteralExpression>());
-    EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->value, 2);
-    EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->suffix,
-              ast::IntLiteralExpression::Suffix::kNone);
-    ASSERT_TRUE(attr_1->Is<ast::StageAttribute>());
-    EXPECT_EQ(attr_1->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
-}
-
 TEST_F(ParserImplTest, AttributeList_Parses) {
     auto p = parser("@workgroup_size(2) @compute");
     auto attrs = p->attribute_list();
@@ -81,14 +54,5 @@
     EXPECT_EQ(p->error(), "1:2: expected attribute");
 }
 
-TEST_F(ParserImplTest, AttributeList_BadAttribute) {
-    auto p = parser("@stage()");
-    auto attrs = p->attribute_list();
-    EXPECT_TRUE(p->has_error());
-    EXPECT_TRUE(attrs.errored);
-    EXPECT_FALSE(attrs.matched);
-    EXPECT_EQ(p->error(), "1:8: invalid value for stage attribute");
-}
-
 }  // namespace
 }  // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
index 60e8f86..186c86a 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
@@ -397,61 +397,6 @@
     EXPECT_EQ(p->error(), "1:21: expected ')' for workgroup_size attribute");
 }
 
-// TODO(crbug.com/tint/1503): Remove when @stage is removed
-TEST_F(ParserImplTest, Attribute_Stage) {
-    auto p = parser("stage(compute)");
-    auto attr = p->attribute();
-    EXPECT_TRUE(attr.matched);
-    EXPECT_FALSE(attr.errored);
-    ASSERT_NE(attr.value, nullptr) << p->error();
-    ASSERT_FALSE(p->has_error());
-    auto* func_attr = attr.value->As<ast::Attribute>();
-    ASSERT_NE(func_attr, nullptr);
-    ASSERT_TRUE(func_attr->Is<ast::StageAttribute>());
-    EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingValue) {
-    auto p = parser("stage()");
-    auto attr = p->attribute();
-    EXPECT_FALSE(attr.matched);
-    EXPECT_TRUE(attr.errored);
-    EXPECT_EQ(attr.value, nullptr);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), "1:7: invalid value for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingInvalid) {
-    auto p = parser("stage(nan)");
-    auto attr = p->attribute();
-    EXPECT_FALSE(attr.matched);
-    EXPECT_TRUE(attr.errored);
-    EXPECT_EQ(attr.value, nullptr);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), "1:7: invalid value for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingLeftParen) {
-    auto p = parser("stage compute)");
-    auto attr = p->attribute();
-    EXPECT_FALSE(attr.matched);
-    EXPECT_TRUE(attr.errored);
-    EXPECT_EQ(attr.value, nullptr);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), "1:7: expected '(' for stage attribute");
-}
-
-TEST_F(ParserImplTest, Attribute_Stage_MissingRightParen) {
-    auto p = parser("stage(compute");
-    auto attr = p->attribute();
-    EXPECT_FALSE(attr.matched);
-    EXPECT_TRUE(attr.errored);
-    EXPECT_EQ(attr.value, nullptr);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), R"(1:1: use of deprecated language feature: remove stage and use @compute
-1:14: expected ')' for stage attribute)");
-}
-
 TEST_F(ParserImplTest, Attribute_Compute) {
     auto p = parser("compute");
     auto attr = p->attribute();
diff --git a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
index 86e4383..c465e4b 100644
--- a/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_constant_decl_test.cc
@@ -24,79 +24,10 @@
     EXPECT_FALSE(attrs.errored);
     EXPECT_FALSE(attrs.matched);
     auto e = p->global_constant_decl(attrs.value);
-    EXPECT_FALSE(p->has_error()) << p->error();
-    EXPECT_TRUE(e.matched);
-    EXPECT_FALSE(e.errored);
-    auto* const_ = e.value->As<ast::Const>();
-    ASSERT_NE(const_, nullptr);
-
-    EXPECT_EQ(const_->symbol, p->builder().Symbols().Get("a"));
-    ASSERT_NE(const_->type, nullptr);
-    EXPECT_TRUE(const_->type->Is<ast::F32>());
-
-    EXPECT_EQ(const_->source.range.begin.line, 1u);
-    EXPECT_EQ(const_->source.range.begin.column, 5u);
-    EXPECT_EQ(const_->source.range.end.line, 1u);
-    EXPECT_EQ(const_->source.range.end.column, 6u);
-
-    ASSERT_NE(const_->initializer, nullptr);
-    EXPECT_TRUE(const_->initializer->Is<ast::LiteralExpression>());
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_Inferred) {
-    auto p = parser("let a = 1.");
-    auto attrs = p->attribute_list();
-    EXPECT_FALSE(attrs.errored);
-    EXPECT_FALSE(attrs.matched);
-    auto e = p->global_constant_decl(attrs.value);
-    EXPECT_FALSE(p->has_error()) << p->error();
-    EXPECT_TRUE(e.matched);
-    EXPECT_FALSE(e.errored);
-    auto* const_ = e.value->As<ast::Const>();
-    ASSERT_NE(const_, nullptr);
-
-    EXPECT_EQ(const_->symbol, p->builder().Symbols().Get("a"));
-    EXPECT_EQ(const_->type, nullptr);
-
-    EXPECT_EQ(const_->source.range.begin.line, 1u);
-    EXPECT_EQ(const_->source.range.begin.column, 5u);
-    EXPECT_EQ(const_->source.range.end.line, 1u);
-    EXPECT_EQ(const_->source.range.end.column, 6u);
-
-    ASSERT_NE(const_->initializer, nullptr);
-    EXPECT_TRUE(const_->initializer->Is<ast::LiteralExpression>());
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_InvalidExpression) {
-    auto p = parser("let a : f32 = if (a) {}");
-    auto attrs = p->attribute_list();
-    EXPECT_FALSE(attrs.errored);
-    EXPECT_FALSE(attrs.matched);
-    auto e = p->global_constant_decl(attrs.value);
     EXPECT_TRUE(p->has_error());
-    EXPECT_TRUE(e.errored);
     EXPECT_FALSE(e.matched);
-    EXPECT_EQ(e.value, nullptr);
-    EXPECT_EQ(
-        p->error(),
-        R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:15: missing initializer for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalLetDecl_MissingExpression) {
-    auto p = parser("let a : f32 =");
-    auto attrs = p->attribute_list();
-    EXPECT_FALSE(attrs.errored);
-    EXPECT_FALSE(attrs.matched);
-    auto e = p->global_constant_decl(attrs.value);
-    EXPECT_TRUE(p->has_error());
     EXPECT_TRUE(e.errored);
-    EXPECT_FALSE(e.matched);
-    EXPECT_EQ(e.value, nullptr);
-    EXPECT_EQ(
-        p->error(),
-        R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:14: missing initializer for 'let' declaration)");
+    EXPECT_EQ(p->error(), "1:1: module-scope 'let' is invalid, use 'const'");
 }
 
 TEST_F(ParserImplTest, GlobalConstDecl) {
diff --git a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
index 6c943c8..2f124d5 100644
--- a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
@@ -58,44 +58,11 @@
 
 TEST_F(ParserImplTest, GlobalDecl_GlobalLet) {
     auto p = parser("let a : i32 = 2;");
-    p->global_decl();
-    ASSERT_FALSE(p->has_error()) << p->error();
-
-    auto program = p->program();
-    ASSERT_EQ(program.AST().GlobalVariables().Length(), 1u);
-
-    auto* v = program.AST().GlobalVariables()[0];
-    EXPECT_EQ(v->symbol, program.Symbols().Get("a"));
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_MissingInitializer) {
-    auto p = parser("let a : vec2<i32>;");
-    p->global_decl();
-    ASSERT_TRUE(p->has_error());
-    EXPECT_EQ(
-        p->error(),
-        R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:18: expected '=' for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_Invalid) {
-    auto p = parser("let a : vec2<i32> 1.0;");
-    p->global_decl();
-    ASSERT_TRUE(p->has_error());
-    EXPECT_EQ(
-        p->error(),
-        R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:19: expected '=' for 'let' declaration)");
-}
-
-TEST_F(ParserImplTest, GlobalDecl_GlobalLet_MissingSemicolon) {
-    auto p = parser("let a : vec2<i32> = vec2<i32>(1, 2)");
-    p->global_decl();
-    ASSERT_TRUE(p->has_error());
-    EXPECT_EQ(
-        p->error(),
-        R"(1:1: use of deprecated language feature: module-scope 'let' has been replaced with 'const'
-1:36: expected ';' for 'const' declaration)");
+    auto e = p->global_decl();
+    EXPECT_TRUE(p->has_error());
+    EXPECT_FALSE(e.matched);
+    EXPECT_TRUE(e.errored);
+    EXPECT_EQ(p->error(), "1:1: module-scope 'let' is invalid, use 'const'");
 }
 
 TEST_F(ParserImplTest, GlobalDecl_GlobalConst) {
diff --git a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc b/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
deleted file mode 100644
index 3c730e2..0000000
--- a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2020 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/reader/wgsl/parser_impl_test_helper.h"
-
-namespace tint::reader::wgsl {
-namespace {
-
-struct PipelineStageData {
-    std::string input;
-    ast::PipelineStage result;
-};
-inline std::ostream& operator<<(std::ostream& out, PipelineStageData data) {
-    return out << data.input;
-}
-
-class PipelineStageTest : public ParserImplTestWithParam<PipelineStageData> {};
-
-TEST_P(PipelineStageTest, Parses) {
-    auto params = GetParam();
-    auto p = parser(params.input);
-
-    auto stage = p->expect_pipeline_stage();
-    ASSERT_FALSE(p->has_error()) << p->error();
-    ASSERT_FALSE(stage.errored);
-    EXPECT_EQ(stage.value, params.result);
-    EXPECT_EQ(stage.source.range.begin.line, 1u);
-    EXPECT_EQ(stage.source.range.begin.column, 1u);
-    EXPECT_EQ(stage.source.range.end.line, 1u);
-    EXPECT_EQ(stage.source.range.end.column, 1u + params.input.size());
-
-    auto& t = p->next();
-    EXPECT_TRUE(t.IsEof());
-}
-INSTANTIATE_TEST_SUITE_P(
-    ParserImplTest,
-    PipelineStageTest,
-    testing::Values(PipelineStageData{"vertex", ast::PipelineStage::kVertex},
-                    PipelineStageData{"fragment", ast::PipelineStage::kFragment},
-                    PipelineStageData{"compute", ast::PipelineStage::kCompute}));
-
-TEST_F(ParserImplTest, PipelineStage_NoMatch) {
-    auto p = parser("not-a-stage");
-    auto stage = p->expect_pipeline_stage();
-    ASSERT_TRUE(p->has_error());
-    ASSERT_TRUE(stage.errored);
-    ASSERT_EQ(p->error(), "1:1: invalid value for stage attribute");
-}
-
-}  // namespace
-}  // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
index 0ad3cc1..8e178ab 100644
--- a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc
@@ -21,74 +21,67 @@
 TEST_P(ParserImplReservedKeywordTest, Function) {
     auto name = GetParam();
     auto p = parser("fn " + name + "() {}");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:4: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:4: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, ModuleConst) {
     auto name = GetParam();
     auto p = parser("const " + name + " : i32 = 1;");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:7: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:7: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, ModuleVar) {
     auto name = GetParam();
     auto p = parser("var " + name + " : i32 = 1;");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:5: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:5: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, FunctionLet) {
     auto name = GetParam();
     auto p = parser("fn f() { let " + name + " : i32 = 1; }");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:14: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:14: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, FunctionVar) {
     auto name = GetParam();
     auto p = parser("fn f() { var " + name + " : i32 = 1; }");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:14: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:14: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, FunctionParam) {
     auto name = GetParam();
     auto p = parser("fn f(" + name + " : i32) {}");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:6: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:6: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, Struct) {
     auto name = GetParam();
     auto p = parser("struct " + name + " {};");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:8: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), R"(1:8: ')" + name + R"(' is a reserved keyword
+1:)" + std::to_string(9 + name.length()) +
+                              R"(: statement found outside of function body)");
 }
 TEST_P(ParserImplReservedKeywordTest, StructMember) {
     auto name = GetParam();
     auto p = parser("struct S { " + name + " : i32, };");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:12: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:12: '" + name + "' is a reserved keyword");
 }
 TEST_P(ParserImplReservedKeywordTest, Alias) {
     auto name = GetParam();
     auto p = parser("type " + name + " = i32;");
-    EXPECT_TRUE(p->Parse());
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(p->error(),
-              "1:6: use of deprecated language feature: '" + name + "' is a reserved keyword");
+    EXPECT_FALSE(p->Parse());
+    EXPECT_TRUE(p->has_error());
+    EXPECT_EQ(p->error(), "1:6: '" + name + "' is a reserved keyword");
 }
 INSTANTIATE_TEST_SUITE_P(ParserImplReservedKeywordTest,
                          ParserImplReservedKeywordTest,
diff --git a/src/tint/resolver/builtin_test.cc b/src/tint/resolver/builtin_test.cc
index c3911fb..3d8a4db 100644
--- a/src/tint/resolver/builtin_test.cc
+++ b/src/tint/resolver/builtin_test.cc
@@ -748,7 +748,7 @@
     EXPECT_EQ(r()->error(), R"(error: no matching call to cross()
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -761,7 +761,7 @@
     EXPECT_EQ(r()->error(), R"(error: no matching call to cross(f32, f32)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -775,7 +775,7 @@
               R"(error: no matching call to cross(vec3<i32>, vec3<i32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -790,7 +790,7 @@
               R"(error: no matching call to cross(vec4<f32>, vec4<f32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -806,7 +806,7 @@
               R"(error: no matching call to cross(vec3<f32>, vec3<f32>, vec3<f32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index afb9f5d..2c012a1 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -37,6 +37,7 @@
 #include "src/tint/sem/type_initializer.h"
 #include "src/tint/sem/u32.h"
 #include "src/tint/sem/vector.h"
+#include "src/tint/utils/bitcast.h"
 #include "src/tint/utils/compiler_macros.h"
 #include "src/tint/utils/map.h"
 #include "src/tint/utils/scoped_assignment.h"
@@ -680,6 +681,33 @@
 }
 
 template <typename NumberT>
+utils::Result<NumberT> ConstEval::Sub(NumberT a, NumberT b) {
+    NumberT result;
+    if constexpr (IsAbstract<NumberT>) {
+        // Check for over/underflow for abstract values
+        if (auto r = CheckedSub(a, b)) {
+            result = r->value;
+        } else {
+            AddError(OverflowErrorMessage(a, "-", b), *current_source);
+            return utils::Failure;
+        }
+    } else {
+        using T = UnwrapNumber<NumberT>;
+        auto sub_values = [](T lhs, T rhs) {
+            if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
+                // Ensure no UB for signed overflow
+                using UT = std::make_unsigned_t<T>;
+                return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
+            } else {
+                return lhs - rhs;
+            }
+        };
+        result = sub_values(a.value, b.value);
+    }
+    return result;
+}
+
+template <typename NumberT>
 utils::Result<NumberT> ConstEval::Mul(NumberT a, NumberT b) {
     using T = UnwrapNumber<NumberT>;
     NumberT result;
@@ -793,6 +821,37 @@
     return r;
 }
 
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2) {
+    auto r1 = Mul(a1, b2);
+    if (!r1) {
+        return utils::Failure;
+    }
+    auto r2 = Mul(b1, a2);
+    if (!r2) {
+        return utils::Failure;
+    }
+    auto r = Sub(r1.Get(), r2.Get());
+    if (!r) {
+        return utils::Failure;
+    }
+    return r;
+}
+
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Clamp(NumberT e, NumberT low, NumberT high) {
+    return NumberT{std::min(std::max(e, low), high)};
+}
+
+auto ConstEval::ClampFunc(const sem::Type* elem_ty) {
+    return [=](auto e, auto low, auto high) -> ImplResult {
+        if (auto r = Clamp(e, low, high)) {
+            return CreateElement(builder, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 auto ConstEval::AddFunc(const sem::Type* elem_ty) {
     return [=](auto a1, auto a2) -> ImplResult {
         if (auto r = Add(a1, a2)) {
@@ -802,6 +861,15 @@
     };
 }
 
+auto ConstEval::SubFunc(const sem::Type* elem_ty) {
+    return [=](auto a1, auto a2) -> ImplResult {
+        if (auto r = Sub(a1, a2)) {
+            return CreateElement(builder, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 auto ConstEval::MulFunc(const sem::Type* elem_ty) {
     return [=](auto a1, auto a2) -> ImplResult {
         if (auto r = Mul(a1, a2)) {
@@ -839,6 +907,15 @@
         };
 }
 
+auto ConstEval::Det2Func(const sem::Type* elem_ty) {
+    return [=](auto a, auto b, auto c, auto d) -> ImplResult {
+        if (auto r = Det2(a, b, c, d)) {
+            return CreateElement(builder, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 ConstEval::Result ConstEval::Literal(const sem::Type* ty, const ast::LiteralExpression* literal) {
     return Switch(
         literal,
@@ -1100,34 +1177,9 @@
 ConstEval::Result ConstEval::OpMinus(const sem::Type* ty,
                                      utils::VectorRef<const sem::Constant*> args,
                                      const Source& source) {
+    TINT_SCOPED_ASSIGNMENT(current_source, &source);
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
-        auto create = [&](auto i, auto j) -> ImplResult {
-            using NumberT = decltype(i);
-            NumberT result;
-            if constexpr (IsAbstract<NumberT>) {
-                // Check for over/underflow for abstract values
-                if (auto r = CheckedSub(i, j)) {
-                    result = r->value;
-                } else {
-                    AddError(OverflowErrorMessage(i, "-", j), source);
-                    return utils::Failure;
-                }
-            } else {
-                using T = UnwrapNumber<NumberT>;
-                auto subtract_values = [](T lhs, T rhs) {
-                    if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
-                        // Ensure no UB for signed underflow
-                        using UT = std::make_unsigned_t<T>;
-                        return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
-                    } else {
-                        return lhs - rhs;
-                    }
-                };
-                result = subtract_values(i.value, j.value);
-            }
-            return CreateElement(builder, c0->Type(), result);
-        };
-        return Dispatch_fia_fiu32_f16(create, c0, c1);
+        return Dispatch_fia_fiu32_f16(SubFunc(c0->Type()), c0, c1);
     };
 
     return TransformBinaryElements(builder, ty, transform, args[0], args[1]);
@@ -1612,7 +1664,8 @@
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i < NumberT(-1.0) || i > NumberT(1.0)) {
-                AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)", source);
+                AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::acos(i.value)));
@@ -1641,7 +1694,8 @@
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i < NumberT(-1.0) || i > NumberT(1.0)) {
-                AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)", source);
+                AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::asin(i.value)));
@@ -1687,7 +1741,8 @@
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i <= NumberT(-1.0) || i >= NumberT(1.0)) {
-                AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)", source);
+                AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::atanh(i.value)));
@@ -1727,11 +1782,7 @@
                                    const Source&) {
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1,
                          const sem::Constant* c2) {
-        auto create = [&](auto e, auto low, auto high) {
-            return CreateElement(builder, c0->Type(),
-                                 decltype(e)(std::min(std::max(e, low), high)));
-        };
-        return Dispatch_fia_fiu32_f16(create, c0, c1, c2);
+        return Dispatch_fia_fiu32_f16(ClampFunc(c0->Type()), c0, c1, c2);
     };
     return TransformElements(builder, ty, transform, args[0], args[1], args[2]);
 }
@@ -1790,6 +1841,52 @@
     return TransformElements(builder, ty, transform, args[0]);
 }
 
+ConstEval::Result ConstEval::cross(const sem::Type* ty,
+                                   utils::VectorRef<const sem::Constant*> args,
+                                   const Source& source) {
+    TINT_SCOPED_ASSIGNMENT(current_source, &source);
+    auto* u = args[0];
+    auto* v = args[1];
+    auto* elem_ty = u->Type()->As<sem::Vector>()->type();
+
+    // cross product of a v3 is the determinant of the 3x3 matrix:
+    //
+    // |i   j   k |
+    // |u0  u1  u2|
+    // |v0  v1  v2|
+    //
+    // |u1 u2|i  - |u0 u2|j + |u0 u1|k
+    // |v1 v2|     |v0 v2|    |v0 v1|
+    //
+    // |u1 u2|i  + |v0 v2|j + |u0 u1|k
+    // |v1 v2|     |u0 u2|    |v0 v1|
+
+    auto* u0 = u->Index(0);
+    auto* u1 = u->Index(1);
+    auto* u2 = u->Index(2);
+    auto* v0 = v->Index(0);
+    auto* v1 = v->Index(1);
+    auto* v2 = v->Index(2);
+
+    // auto x = Dispatch_fa_f32_f16(ab_minus_cd_func(elem_ty), u->Index(1), v->Index(2),
+    //                              v->Index(1), u->Index(2));
+    auto x = Dispatch_fa_f32_f16(Det2Func(elem_ty), u1, u2, v1, v2);
+    if (!x) {
+        return utils::Failure;
+    }
+    auto y = Dispatch_fa_f32_f16(Det2Func(elem_ty), v0, v2, u0, u2);
+    if (!y) {
+        return utils::Failure;
+    }
+    auto z = Dispatch_fa_f32_f16(Det2Func(elem_ty), u0, u1, v0, v1);
+    if (!z) {
+        return utils::Failure;
+    }
+
+    return CreateComposite(builder, ty,
+                           utils::Vector<const sem::Constant*, 3>{x.Get(), y.Get(), z.Get()});
+}
+
 ConstEval::Result ConstEval::extractBits(const sem::Type* ty,
                                          utils::VectorRef<const sem::Constant*> args,
                                          const Source& source) {
@@ -1804,18 +1901,18 @@
             NumberUT in_offset = args[1]->As<NumberUT>();
             NumberUT in_count = args[2]->As<NumberUT>();
 
-            constexpr UT w = sizeof(UT) * 8;
-            if ((in_offset + in_count) > w) {
-                AddError("'offset + 'count' must be less than or equal to the bit width of 'e'",
-                         source);
-                return utils::Failure;
-            }
-
             // Cast all to unsigned
             UT e = static_cast<UT>(in_e);
             UT o = static_cast<UT>(in_offset);
             UT c = static_cast<UT>(in_count);
 
+            constexpr UT w = sizeof(UT) * 8;
+            if (o > w || c > w || (o + c) > w) {
+                AddError("'offset + 'count' must be less than or equal to the bit width of 'e'",
+                         source);
+                return utils::Failure;
+            }
+
             NumberT result;
             if (c == UT{0}) {
                 // The result is 0 if c is 0
@@ -1979,6 +2076,78 @@
     return TransformElements(builder, ty, transform, args[0], args[1]);
 }
 
+ConstEval::Result ConstEval::pack2x16snorm(const sem::Type* ty,
+                                           utils::VectorRef<const sem::Constant*> args,
+                                           const Source&) {
+    auto calc = [&](f32 val) -> u32 {
+        auto clamped = Clamp(val, f32(-1.0f), f32(1.0f)).Get();
+        return u32(utils::Bitcast<uint16_t>(
+            static_cast<int16_t>(std::floor(0.5f + (32767.0f * clamped)))));
+    };
+
+    auto* e = args[0];
+    auto e0 = calc(e->Index(0)->As<f32>());
+    auto e1 = calc(e->Index(1)->As<f32>());
+
+    u32 ret = u32((e0 & 0x0000'ffff) | (e1 << 16));
+    return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack2x16unorm(const sem::Type* ty,
+                                           utils::VectorRef<const sem::Constant*> args,
+                                           const Source&) {
+    auto calc = [&](f32 val) -> u32 {
+        auto clamped = Clamp(val, f32(0.0f), f32(1.0f)).Get();
+        return u32{std::floor(0.5f + (65535.0f * clamped))};
+    };
+
+    auto* e = args[0];
+    auto e0 = calc(e->Index(0)->As<f32>());
+    auto e1 = calc(e->Index(1)->As<f32>());
+
+    u32 ret = u32((e0 & 0x0000'ffff) | (e1 << 16));
+    return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack4x8snorm(const sem::Type* ty,
+                                          utils::VectorRef<const sem::Constant*> args,
+                                          const Source&) {
+    auto calc = [&](f32 val) -> u32 {
+        auto clamped = Clamp(val, f32(-1.0f), f32(1.0f)).Get();
+        return u32(
+            utils::Bitcast<uint8_t>(static_cast<int8_t>(std::floor(0.5f + (127.0f * clamped)))));
+    };
+
+    auto* e = args[0];
+    auto e0 = calc(e->Index(0)->As<f32>());
+    auto e1 = calc(e->Index(1)->As<f32>());
+    auto e2 = calc(e->Index(2)->As<f32>());
+    auto e3 = calc(e->Index(3)->As<f32>());
+
+    uint32_t mask = 0x0000'00ff;
+    u32 ret = u32((e0 & mask) | ((e1 & mask) << 8) | ((e2 & mask) << 16) | ((e3 & mask) << 24));
+    return CreateElement(builder, ty, ret);
+}
+
+ConstEval::Result ConstEval::pack4x8unorm(const sem::Type* ty,
+                                          utils::VectorRef<const sem::Constant*> args,
+                                          const Source&) {
+    auto calc = [&](f32 val) -> u32 {
+        auto clamped = Clamp(val, f32(0.0f), f32(1.0f)).Get();
+        return u32{std::floor(0.5f + (255.0f * clamped))};
+    };
+
+    auto* e = args[0];
+    auto e0 = calc(e->Index(0)->As<f32>());
+    auto e1 = calc(e->Index(1)->As<f32>());
+    auto e2 = calc(e->Index(2)->As<f32>());
+    auto e3 = calc(e->Index(3)->As<f32>());
+
+    uint32_t mask = 0x0000'00ff;
+    u32 ret = u32((e0 & mask) | ((e1 & mask) << 8) | ((e2 & mask) << 16) | ((e3 & mask) << 24));
+    return CreateElement(builder, ty, ret);
+}
+
 ConstEval::Result ConstEval::reverseBits(const sem::Type* ty,
                                          utils::VectorRef<const sem::Constant*> args,
                                          const Source&) {
@@ -2085,6 +2254,68 @@
     return TransformElements(builder, ty, transform, args[0], args[1]);
 }
 
+ConstEval::Result ConstEval::unpack2x16snorm(const sem::Type* ty,
+                                             utils::VectorRef<const sem::Constant*> args,
+                                             const Source&) {
+    auto* inner_ty = sem::Type::DeepestElementOf(ty);
+    auto e = args[0]->As<u32>().value;
+
+    utils::Vector<const sem::Constant*, 2> els;
+    els.Reserve(2);
+    for (size_t i = 0; i < 2; ++i) {
+        auto val = f32(
+            std::max(static_cast<float>(int16_t((e >> (16 * i)) & 0x0000'ffff)) / 32767.f, -1.f));
+        els.Push(CreateElement(builder, inner_ty, val));
+    }
+    return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack2x16unorm(const sem::Type* ty,
+                                             utils::VectorRef<const sem::Constant*> args,
+                                             const Source&) {
+    auto* inner_ty = sem::Type::DeepestElementOf(ty);
+    auto e = args[0]->As<u32>().value;
+
+    utils::Vector<const sem::Constant*, 2> els;
+    els.Reserve(2);
+    for (size_t i = 0; i < 2; ++i) {
+        auto val = f32(static_cast<float>(uint16_t((e >> (16 * i)) & 0x0000'ffff)) / 65535.f);
+        els.Push(CreateElement(builder, inner_ty, val));
+    }
+    return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack4x8snorm(const sem::Type* ty,
+                                            utils::VectorRef<const sem::Constant*> args,
+                                            const Source&) {
+    auto* inner_ty = sem::Type::DeepestElementOf(ty);
+    auto e = args[0]->As<u32>().value;
+
+    utils::Vector<const sem::Constant*, 4> els;
+    els.Reserve(4);
+    for (size_t i = 0; i < 4; ++i) {
+        auto val =
+            f32(std::max(static_cast<float>(int8_t((e >> (8 * i)) & 0x0000'00ff)) / 127.f, -1.f));
+        els.Push(CreateElement(builder, inner_ty, val));
+    }
+    return CreateComposite(builder, ty, std::move(els));
+}
+
+ConstEval::Result ConstEval::unpack4x8unorm(const sem::Type* ty,
+                                            utils::VectorRef<const sem::Constant*> args,
+                                            const Source&) {
+    auto* inner_ty = sem::Type::DeepestElementOf(ty);
+    auto e = args[0]->As<u32>().value;
+
+    utils::Vector<const sem::Constant*, 4> els;
+    els.Reserve(4);
+    for (size_t i = 0; i < 4; ++i) {
+        auto val = f32(static_cast<float>(uint8_t((e >> (8 * i)) & 0x0000'00ff)) / 255.f);
+        els.Push(CreateElement(builder, inner_ty, val));
+    }
+    return CreateComposite(builder, ty, std::move(els));
+}
+
 ConstEval::Result ConstEval::quantizeToF16(const sem::Type* ty,
                                            utils::VectorRef<const sem::Constant*> args,
                                            const Source&) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index f309150..9907b36 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -503,6 +503,15 @@
                               utils::VectorRef<const sem::Constant*> args,
                               const Source& source);
 
+    /// cross builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result cross(const sem::Type* ty,
+                 utils::VectorRef<const sem::Constant*> args,
+                 const Source& source);
+
     /// extractBits builtin
     /// @param ty the expression type
     /// @param args the input arguments
@@ -548,6 +557,42 @@
                       utils::VectorRef<const sem::Constant*> args,
                       const Source& source);
 
+    /// pack2x16snorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result pack2x16snorm(const sem::Type* ty,
+                         utils::VectorRef<const sem::Constant*> args,
+                         const Source& source);
+
+    /// pack2x16unorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result pack2x16unorm(const sem::Type* ty,
+                         utils::VectorRef<const sem::Constant*> args,
+                         const Source& source);
+
+    /// pack4x8snorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result pack4x8snorm(const sem::Type* ty,
+                        utils::VectorRef<const sem::Constant*> args,
+                        const Source& source);
+
+    /// pack4x8unorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result pack4x8unorm(const sem::Type* ty,
+                        utils::VectorRef<const sem::Constant*> args,
+                        const Source& source);
+
     /// reverseBits builtin
     /// @param ty the expression type
     /// @param args the input arguments
@@ -602,6 +647,42 @@
                 utils::VectorRef<const sem::Constant*> args,
                 const Source& source);
 
+    /// unpack2x16snorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result unpack2x16snorm(const sem::Type* ty,
+                           utils::VectorRef<const sem::Constant*> args,
+                           const Source& source);
+
+    /// unpack2x16unorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result unpack2x16unorm(const sem::Type* ty,
+                           utils::VectorRef<const sem::Constant*> args,
+                           const Source& source);
+
+    /// unpack4x8snorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result unpack4x8snorm(const sem::Type* ty,
+                          utils::VectorRef<const sem::Constant*> args,
+                          const Source& source);
+
+    /// unpack4x8unorm builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result unpack4x8unorm(const sem::Type* ty,
+                          utils::VectorRef<const sem::Constant*> args,
+                          const Source& source);
+
     /// quantizeToF16 builtin
     /// @param ty the expression type
     /// @param args the input arguments
@@ -625,6 +706,13 @@
     template <typename NumberT>
     utils::Result<NumberT> Add(NumberT a, NumberT b);
 
+    /// Subtracts two Number<T>s
+    /// @param a the lhs number
+    /// @param b the rhs number
+    /// @returns the result number on success, or logs an error and returns Failure
+    template <typename NumberT>
+    utils::Result<NumberT> Sub(NumberT a, NumberT b);
+
     /// Multiplies two Number<T>s
     /// @param a the lhs number
     /// @param b the rhs number
@@ -677,12 +765,34 @@
                                 NumberT b3,
                                 NumberT b4);
 
+    /// Returns the determinant of the 2x2 matrix [(a1, a2), (b1, b2)]
+    /// @param a1 component 1 of the first column vector
+    /// @param a2 component 2 of the first column vector
+    /// @param b1 component 1 of the second column vector
+    /// @param b2 component 2 of the second column vector
+    template <typename NumberT>
+    utils::Result<NumberT> Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2);
+
+    /// Clamps e between low and high
+    /// @param e the number to clamp
+    /// @param low the lower bound
+    /// @param high the upper bound
+    /// @returns the result number on success, or logs an error and returns Failure
+    template <typename NumberT>
+    utils::Result<NumberT> Clamp(NumberT e, NumberT low, NumberT high);
+
     /// Returns a callable that calls Add, and creates a Constant with its result of type `elem_ty`
     /// if successful, or returns Failure otherwise.
     /// @param elem_ty the element type of the Constant to create on success
     /// @returns the callable function
     auto AddFunc(const sem::Type* elem_ty);
 
+    /// Returns a callable that calls Sub, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto SubFunc(const sem::Type* elem_ty);
+
     /// Returns a callable that calls Mul, and creates a Constant with its result of type `elem_ty`
     /// if successful, or returns Failure otherwise.
     /// @param elem_ty the element type of the Constant to create on success
@@ -707,6 +817,18 @@
     /// @returns the callable function
     auto Dot4Func(const sem::Type* elem_ty);
 
+    /// Returns a callable that calls Det2, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto Det2Func(const sem::Type* elem_ty);
+
+    /// Returns a callable that calls Clamp, and creates a Constant with its result of type
+    /// `elem_ty` if successful, or returns Failure otherwise.
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto ClampFunc(const sem::Type* elem_ty);
+
     ProgramBuilder& builder;
     const Source* current_source = nullptr;
 };
diff --git a/src/tint/resolver/const_eval_builtin_test.cc b/src/tint/resolver/const_eval_builtin_test.cc
index 3dc395b..fddca38 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -9,7 +9,7 @@
 // Unless required by applicable law or agreed to in writing, software
 // distributed under the License is distributed on an "AS IS" BASIS,
 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
+// See the License for the empecific language governing permissions and
 // limitations under the License.
 
 #include "src/tint/resolver/const_eval_test.h"
@@ -17,6 +17,7 @@
 #include "src/tint/utils/result.h"
 
 using namespace tint::number_suffixes;  // NOLINT
+using ::testing::HasSubstr;
 
 namespace tint::resolver {
 namespace {
@@ -28,8 +29,8 @@
     Case(utils::VectorRef<Types> in_args, Types expected_value)
         : args(std::move(in_args)), expected(Success{std::move(expected_value), false, false}) {}
 
-    Case(utils::VectorRef<Types> in_args, const char* expected_err)
-        : args(std::move(in_args)), expected(Failure{expected_err}) {}
+    Case(utils::VectorRef<Types> in_args, std::string expected_err)
+        : args(std::move(in_args)), expected(Failure{std::move(expected_err)}) {}
 
     /// Expected value may be positive or negative
     Case& PosOrNeg() {
@@ -53,7 +54,7 @@
         bool float_compare = false;
     };
     struct Failure {
-        const char* error = nullptr;
+        std::string error = nullptr;
     };
 
     utils::Vector<Types, 8> args;
@@ -94,17 +95,27 @@
 }
 
 /// Creates a Case with Values for args and expected error
-static Case E(std::initializer_list<Types> args, const char* err) {
-    return Case{utils::Vector<Types, 8>{args}, err};
+static Case E(std::initializer_list<Types> args, std::string err) {
+    return Case{utils::Vector<Types, 8>{args}, std::move(err)};
 }
 
 /// Convenience overload that creates an expected-error Case with just scalars
-static Case E(std::initializer_list<ScalarTypes> sargs, const char* err) {
+static Case E(std::initializer_list<ScalarTypes> sargs, std::string err) {
     utils::Vector<Types, 8> args;
     for (auto& sa : sargs) {
         std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa);
     }
-    return Case{std::move(args), err};
+    return Case{std::move(args), std::move(err)};
+}
+
+/// Returns the overflow error message for binary ops
+template <typename NumberT>
+std::string OverflowErrorMessage(NumberT lhs, const char* op, NumberT rhs) {
+    std::stringstream ss;
+    ss << std::setprecision(20);
+    ss << "'" << lhs.value << " " << op << " " << rhs.value << "' cannot be represented as '"
+       << FriendlyName<NumberT>() << "'";
+    return ss.str();
 }
 
 using ResolverConstEvalBuiltinTest = ResolverTestWithParam<std::tuple<sem::BuiltinType, Case>>;
@@ -776,6 +787,130 @@
                      testing::ValuesIn(Concat(CountOneBitsCases<i32>(),  //
                                               CountOneBitsCases<u32>()))));
 
+template <typename T, bool finite_only>
+std::vector<Case> CrossCases() {
+    constexpr auto vec_x = [](T v) { return Vec(T(v), T(0), T(0)); };
+    constexpr auto vec_y = [](T v) { return Vec(T(0), T(v), T(0)); };
+    constexpr auto vec_z = [](T v) { return Vec(T(0), T(0), T(v)); };
+
+    const auto zero = Vec(T(0), T(0), T(0));
+    const auto unit_x = vec_x(T(1));
+    const auto unit_y = vec_y(T(1));
+    const auto unit_z = vec_z(T(1));
+    const auto neg_unit_x = vec_x(-T(1));
+    const auto neg_unit_y = vec_y(-T(1));
+    const auto neg_unit_z = vec_z(-T(1));
+    const auto highest_x = vec_x(T::Highest());
+    const auto highest_y = vec_y(T::Highest());
+    const auto highest_z = vec_z(T::Highest());
+    const auto smallest_x = vec_x(T::Smallest());
+    const auto smallest_y = vec_y(T::Smallest());
+    const auto smallest_z = vec_z(T::Smallest());
+    const auto lowest_x = vec_x(T::Lowest());
+    const auto lowest_y = vec_y(T::Lowest());
+    const auto lowest_z = vec_z(T::Lowest());
+    const auto inf_x = vec_x(T::Inf());
+    const auto inf_y = vec_y(T::Inf());
+    const auto inf_z = vec_z(T::Inf());
+    const auto neg_inf_x = vec_x(-T::Inf());
+    const auto neg_inf_y = vec_y(-T::Inf());
+    const auto neg_inf_z = vec_z(-T::Inf());
+
+    std::vector<Case> r = {
+        C({zero, zero}, zero),
+
+        C({unit_x, unit_x}, zero),
+        C({unit_y, unit_y}, zero),
+        C({unit_z, unit_z}, zero),
+
+        C({smallest_x, smallest_x}, zero),
+        C({smallest_y, smallest_y}, zero),
+        C({smallest_z, smallest_z}, zero),
+
+        C({lowest_x, lowest_x}, zero),
+        C({lowest_y, lowest_y}, zero),
+        C({lowest_z, lowest_z}, zero),
+
+        C({highest_x, highest_x}, zero),
+        C({highest_y, highest_y}, zero),
+        C({highest_z, highest_z}, zero),
+
+        C({smallest_x, highest_x}, zero),
+        C({smallest_y, highest_y}, zero),
+        C({smallest_z, highest_z}, zero),
+
+        C({unit_x, neg_unit_x}, zero).PosOrNeg(),
+        C({unit_y, neg_unit_y}, zero).PosOrNeg(),
+        C({unit_z, neg_unit_z}, zero).PosOrNeg(),
+
+        C({unit_x, unit_y}, unit_z),
+        C({unit_y, unit_x}, neg_unit_z),
+
+        C({unit_z, unit_x}, unit_y),
+        C({unit_x, unit_z}, neg_unit_y),
+
+        C({unit_y, unit_z}, unit_x),
+        C({unit_z, unit_y}, neg_unit_x),
+
+        C({vec_x(T(1)), vec_y(T(2))}, vec_z(T(2))),
+        C({vec_y(T(1)), vec_x(T(2))}, vec_z(-T(2))),
+        C({vec_x(T(2)), vec_y(T(3))}, vec_z(T(6))),
+        C({vec_y(T(2)), vec_x(T(3))}, vec_z(-T(6))),
+
+        C({Vec(T(1), T(2), T(3)), Vec(T(1), T(5), T(7))}, Vec(T(-1), T(-4), T(3))),
+        C({Vec(T(33), T(44), T(55)), Vec(T(13), T(42), T(39))}, Vec(T(-594), T(-572), T(814))),
+        C({Vec(T(3.5), T(4), T(5.5)), Vec(T(1), T(4.5), T(3.5))},
+          Vec(T(-10.75), T(-6.75), T(11.75))),
+    };
+
+    ConcatIntoIf<!finite_only>(  //
+        r, std::vector<Case>{
+               C({highest_x, highest_y}, inf_z).PosOrNeg(),  //
+               C({highest_y, highest_x}, inf_z).PosOrNeg(),  //
+               C({highest_z, highest_x}, inf_y).PosOrNeg(),  //
+               C({highest_x, highest_z}, inf_y).PosOrNeg(),  //
+               C({highest_y, highest_z}, inf_x).PosOrNeg(),  //
+               C({highest_z, highest_y}, inf_x).PosOrNeg(),  //
+               C({lowest_x, lowest_y}, inf_z).PosOrNeg(),    //
+               C({lowest_y, lowest_x}, inf_z).PosOrNeg(),    //
+               C({lowest_z, lowest_x}, inf_y).PosOrNeg(),    //
+               C({lowest_x, lowest_z}, inf_y).PosOrNeg(),    //
+               C({lowest_y, lowest_z}, inf_x).PosOrNeg(),    //
+               C({lowest_z, lowest_y}, inf_x).PosOrNeg(),
+           });
+
+    std::string pos_error_msg =
+        "12:34 error: " + OverflowErrorMessage(T::Highest(), "*", T::Highest());
+    std::string neg_error_msg =
+        "12:34 error: " + OverflowErrorMessage(T::Lowest(), "*", T::Lowest());
+
+    ConcatIntoIf<finite_only>(  //
+        r, std::vector<Case>{
+               E({highest_x, highest_y}, pos_error_msg),
+               E({highest_y, highest_x}, pos_error_msg),
+               E({highest_z, highest_x}, pos_error_msg),
+               E({highest_x, highest_z}, pos_error_msg),
+               E({highest_y, highest_z}, pos_error_msg),
+               E({highest_z, highest_y}, pos_error_msg),
+               E({lowest_x, lowest_y}, neg_error_msg),
+               E({lowest_y, lowest_x}, neg_error_msg),
+               E({lowest_z, lowest_x}, neg_error_msg),
+               E({lowest_x, lowest_z}, neg_error_msg),
+               E({lowest_y, lowest_z}, neg_error_msg),
+               E({lowest_z, lowest_y}, neg_error_msg),
+           });
+
+    return r;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Cross,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kCross),
+                     testing::ValuesIn(Concat(CrossCases<AFloat, true>(),  //
+                                              CrossCases<f32, false>(),
+                                              CrossCases<f32, false>(),  //
+                                              CrossCases<f16, false>()))));
+
 template <typename T>
 std::vector<Case> FirstLeadingBitCases() {
     using B = BitValues<T>;
@@ -1091,8 +1226,77 @@
                              std::make_tuple(33, 33),             //
                              std::make_tuple(34, 34),             //
                              std::make_tuple(1000, 1000),         //
+                             std::make_tuple(u32::Highest(), 1),  //
+                             std::make_tuple(1, u32::Highest()),  //
                              std::make_tuple(u32::Highest(), u32::Highest())));
 
+std::vector<Case> Pack4x8snormCases() {
+    return {
+        C({Vec(f32(0), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0000))),
+        C({Vec(f32(0), f32(0), f32(0), f32(-1))}, Val(u32(0x8100'0000))),
+        C({Vec(f32(0), f32(0), f32(0), f32(1))}, Val(u32(0x7f00'0000))),
+        C({Vec(f32(0), f32(0), f32(-1), f32(0))}, Val(u32(0x0081'0000))),
+        C({Vec(f32(0), f32(1), f32(0), f32(0))}, Val(u32(0x0000'7f00))),
+        C({Vec(f32(-1), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0081))),
+        C({Vec(f32(1), f32(-1), f32(1), f32(-1))}, Val(u32(0x817f'817f))),
+        C({Vec(f32::Highest(), f32(-0.5), f32(0.5), f32::Lowest())}, Val(u32(0x8140'c17f))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Pack4x8snorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kPack4X8Snorm),
+                     testing::ValuesIn(Pack4x8snormCases())));
+
+std::vector<Case> Pack4x8unormCases() {
+    return {
+        C({Vec(f32(0), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0000))),
+        C({Vec(f32(0), f32(0), f32(0), f32(1))}, Val(u32(0xff00'0000))),
+        C({Vec(f32(0), f32(0), f32(1), f32(0))}, Val(u32(0x00ff'0000))),
+        C({Vec(f32(0), f32(1), f32(0), f32(0))}, Val(u32(0x0000'ff00))),
+        C({Vec(f32(1), f32(0), f32(0), f32(0))}, Val(u32(0x0000'00ff))),
+        C({Vec(f32(1), f32(0), f32(1), f32(0))}, Val(u32(0x00ff'00ff))),
+        C({Vec(f32::Highest(), f32(0), f32(0.5), f32::Lowest())}, Val(u32(0x0080'00ff))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Pack4x8unorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kPack4X8Unorm),
+                     testing::ValuesIn(Pack4x8unormCases())));
+
+std::vector<Case> Pack2x16snormCases() {
+    return {
+        C({Vec(f32(0), f32(0))}, Val(u32(0x0000'0000))),
+        C({Vec(f32(0), f32(-1))}, Val(u32(0x8001'0000))),
+        C({Vec(f32(0), f32(1))}, Val(u32(0x7fff'0000))),
+        C({Vec(f32(-1), f32(0))}, Val(u32(0x0000'8001))),
+        C({Vec(f32(1), f32(0))}, Val(u32(0x0000'7fff))),
+        C({Vec(f32(1), f32(-1))}, Val(u32(0x8001'7fff))),
+        C({Vec(f32::Highest(), f32::Lowest())}, Val(u32(0x8001'7fff))),
+        C({Vec(f32(-0.5), f32(0.5))}, Val(u32(0x4000'c001))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Pack2x16snorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kPack2X16Snorm),
+                     testing::ValuesIn(Pack2x16snormCases())));
+
+std::vector<Case> Pack2x16unormCases() {
+    return {
+        C({Vec(f32(0), f32(1))}, Val(u32(0xffff'0000))),
+        C({Vec(f32(1), f32(0))}, Val(u32(0x0000'ffff))),
+        C({Vec(f32(0.5), f32(0))}, Val(u32(0x0000'8000))),
+        C({Vec(f32::Highest(), f32::Lowest())}, Val(u32(0x0000'ffff))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Pack2x16unorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kPack2X16Unorm),
+                     testing::ValuesIn(Pack2x16unormCases())));
+
 template <typename T>
 std::vector<Case> ReverseBitsCases() {
     using B = BitValues<T>;
@@ -1268,6 +1472,74 @@
                                               StepCases<f32>(),
                                               StepCases<f16>()))));
 
+std::vector<Case> Unpack4x8snormCases() {
+    return {
+        C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0), f32(0), f32(0))),
+        C({Val(u32(0x8100'0000))}, Vec(f32(0), f32(0), f32(0), f32(-1))),
+        C({Val(u32(0x7f00'0000))}, Vec(f32(0), f32(0), f32(0), f32(1))),
+        C({Val(u32(0x0081'0000))}, Vec(f32(0), f32(0), f32(-1), f32(0))),
+        C({Val(u32(0x0000'7f00))}, Vec(f32(0), f32(1), f32(0), f32(0))),
+        C({Val(u32(0x0000'0081))}, Vec(f32(-1), f32(0), f32(0), f32(0))),
+        C({Val(u32(0x817f'817f))}, Vec(f32(1), f32(-1), f32(1), f32(-1))),
+        C({Val(u32(0x816d'937f))},
+          Vec(f32(1), f32(-0.8582677165354), f32(0.8582677165354), f32(-1))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Unpack4x8snorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kUnpack4X8Snorm),
+                     testing::ValuesIn(Unpack4x8snormCases())));
+
+std::vector<Case> Unpack4x8unormCases() {
+    return {
+        C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0), f32(0), f32(0))),
+        C({Val(u32(0xff00'0000))}, Vec(f32(0), f32(0), f32(0), f32(1))),
+        C({Val(u32(0x00ff'0000))}, Vec(f32(0), f32(0), f32(1), f32(0))),
+        C({Val(u32(0x0000'ff00))}, Vec(f32(0), f32(1), f32(0), f32(0))),
+        C({Val(u32(0x0000'00ff))}, Vec(f32(1), f32(0), f32(0), f32(0))),
+        C({Val(u32(0x00ff'00ff))}, Vec(f32(1), f32(0), f32(1), f32(0))),
+        C({Val(u32(0x0066'00ff))}, Vec(f32(1), f32(0), f32(0.4), f32(0))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Unpack4x8unorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kUnpack4X8Unorm),
+                     testing::ValuesIn(Unpack4x8unormCases())));
+
+std::vector<Case> Unpack2x16snormCases() {
+    return {
+        C({Val(u32(0x0000'0000))}, Vec(f32(0), f32(0))),
+        C({Val(u32(0x8001'0000))}, Vec(f32(0), f32(-1))),
+        C({Val(u32(0x7fff'0000))}, Vec(f32(0), f32(1))),
+        C({Val(u32(0x0000'8001))}, Vec(f32(-1), f32(0))),
+        C({Val(u32(0x0000'7fff))}, Vec(f32(1), f32(0))),
+        C({Val(u32(0x8001'7fff))}, Vec(f32(1), f32(-1))),
+        C({Val(u32(0x8001'7fff))}, Vec(f32(1), f32(-1))),
+        C({Val(u32(0x4000'999a))}, Vec(f32(-0.80001220740379), f32(0.500015259254737))).FloatComp(),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Unpack2x16snorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kUnpack2X16Snorm),
+                     testing::ValuesIn(Unpack2x16snormCases())));
+
+std::vector<Case> Unpack2x16unormCases() {
+    return {
+        C({Val(u32(0xffff'0000))}, Vec(f32(0), f32(1))),
+        C({Val(u32(0x0000'ffff))}, Vec(f32(1), f32(0))),
+        C({Val(u32(0x0000'6666))}, Vec(f32(0.4), f32(0))),
+        C({Val(u32(0x0000'ffff))}, Vec(f32(1), f32(0))),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Unpack2x16unorm,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kUnpack2X16Unorm),
+                     testing::ValuesIn(Unpack2x16unormCases())));
+
 std::vector<Case> QuantizeToF16Cases() {
     (void)E({Vec(0_f, 0_f)}, "");  // Currently unused, but will be soon.
     return {
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 5a21c55..3238c8b 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -193,8 +193,6 @@
             },
             [&](const ast::Function* func) {
                 Declare(func->symbol, func);
-                TraverseAttributes(func->attributes);
-                TraverseAttributes(func->return_type_attributes);
                 TraverseFunction(func);
             },
             [&](const ast::Variable* var) {
@@ -216,10 +214,13 @@
     /// Traverses the function, performing symbol resolution and determining
     /// global dependencies.
     void TraverseFunction(const ast::Function* func) {
+        TraverseAttributes(func->attributes);
+        TraverseAttributes(func->return_type_attributes);
         // Perform symbol resolution on all the parameter types before registering
         // the parameters themselves. This allows the case of declaring a parameter
         // with the same identifier as its type.
         for (auto* param : func->params) {
+            TraverseAttributes(param->attributes);
             TraverseType(param->type);
         }
         // Resolve the return type
diff --git a/src/tint/resolver/dependency_graph_test.cc b/src/tint/resolver/dependency_graph_test.cc
index 311c9cc..272357b 100644
--- a/src/tint/resolver/dependency_graph_test.cc
+++ b/src/tint/resolver/dependency_graph_test.cc
@@ -1237,9 +1237,14 @@
                                    })});
     GlobalVar(Sym(), T, V);
     GlobalConst(Sym(), T, V);
-    Func(Sym(),                           //
-         utils::Vector{Param(Sym(), T)},  //
-         T,                               // Return type
+    Func(Sym(),
+         utils::Vector{
+             Param(Sym(), T,
+                   utils::Vector{
+                       Location(V),  // Parameter attributes
+                   }),
+         },
+         T,  // Return type
          utils::Vector{
              Decl(Var(Sym(), T, V)),                    //
              Decl(Let(Sym(), T, V)),                    //
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index ddce1ba..dcba75b 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -13766,7 +13766,7 @@
     /* parameters */ &kParameters[841],
     /* return matcher indices */ &kMatcherIndices[102],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::unpack4x8unorm,
   },
   {
     /* [453] */
@@ -13778,7 +13778,7 @@
     /* parameters */ &kParameters[842],
     /* return matcher indices */ &kMatcherIndices[102],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::unpack4x8snorm,
   },
   {
     /* [454] */
@@ -13790,7 +13790,7 @@
     /* parameters */ &kParameters[843],
     /* return matcher indices */ &kMatcherIndices[132],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::unpack2x16unorm,
   },
   {
     /* [455] */
@@ -13802,7 +13802,7 @@
     /* parameters */ &kParameters[844],
     /* return matcher indices */ &kMatcherIndices[132],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::unpack2x16snorm,
   },
   {
     /* [456] */
@@ -13862,7 +13862,7 @@
     /* parameters */ &kParameters[881],
     /* return matcher indices */ &kMatcherIndices[95],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::pack4x8unorm,
   },
   {
     /* [461] */
@@ -13874,7 +13874,7 @@
     /* parameters */ &kParameters[882],
     /* return matcher indices */ &kMatcherIndices[95],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::pack4x8snorm,
   },
   {
     /* [462] */
@@ -13886,7 +13886,7 @@
     /* parameters */ &kParameters[884],
     /* return matcher indices */ &kMatcherIndices[95],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::pack2x16snorm,
   },
   {
     /* [463] */
@@ -13917,12 +13917,12 @@
     /* num parameters */ 2,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[22],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[777],
     /* return matcher indices */ &kMatcherIndices[104],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::cross,
   },
   {
     /* [466] */
@@ -14006,7 +14006,7 @@
     /* parameters */ &kParameters[883],
     /* return matcher indices */ &kMatcherIndices[95],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::pack2x16unorm,
   },
 };
 
@@ -14020,8 +14020,8 @@
   },
   {
     /* [1] */
-    /* fn acos<T : fa_f32_f16>(@test_value(0.87758256189) T) -> T */
-    /* fn acos<N : num, T : fa_f32_f16>(@test_value(0.87758256189) vec<N, T>) -> vec<N, T> */
+    /* fn acos<T : fa_f32_f16>(@test_value(0.96891242171) T) -> T */
+    /* fn acos<N : num, T : fa_f32_f16>(@test_value(0.96891242171) vec<N, T>) -> vec<N, T> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[428],
   },
@@ -14138,7 +14138,7 @@
   },
   {
     /* [18] */
-    /* fn cross<T : f32_f16>(vec3<T>, vec3<T>) -> vec3<T> */
+    /* fn cross<T : fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T> */
     /* num overloads */ 1,
     /* overloads */ &kOverloads[465],
   },
diff --git a/src/tint/tint_gdb.py b/src/tint/tint_gdb.py
index 2dce446..88c614d 100644
--- a/src/tint/tint_gdb.py
+++ b/src/tint/tint_gdb.py
@@ -165,11 +165,11 @@
     'UtilsVector', '^tint::utils::VectorRef<.*>$', UtilsVectorRefPrinter)
 
 
-class UtilsHashsetPrinter(Printer):
-    '''Printer for Hashset<T, N, HASH, EQUAL>'''
+class UtilsHashmapBasePrinter(Printer):
+    '''Base Printer for HashmapBase-derived types'''
 
     def __init__(self, val):
-        super(UtilsHashsetPrinter, self).__init__(val)
+        super(UtilsHashmapBasePrinter, self).__init__(val)
         self.slice = UtilsVectorPrinter(self.val['slots_']).slice_printer()
         self.try_read_std_optional_func = self.try_read_std_optional
 
@@ -185,26 +185,32 @@
         for slot in range(0, self.slice.length()):
             v = self.slice.value_at(slot)
             if v['hash'] != 0:
-                value = v['value']
+                entry = v['entry']
 
-                # value is a std::optional, let's try to extract its value for display
-                kvp = self.try_read_std_optional_func(slot, value)
+                # entry is a std::optional, let's try to extract its value for display
+                kvp = self.try_read_std_optional_func(slot, entry)
                 if kvp is None:
-                    # If we failed, just output the slot and value as is, which will use
+                    # If we failed, just output the slot and entry as is, which will use
                     # the default visualizer for each.
-                    kvp = slot, value
+                    kvp = slot, entry
 
                 yield str(kvp[0]), kvp[1]
 
     def display_hint(self):
         return 'array'
 
-    def try_read_std_optional(self, slot, value):
+    def try_read_std_optional(self, slot, entry):
+        return None
+
+
+class UtilsHashsetPrinter(UtilsHashmapBasePrinter):
+    '''Printer for Hashset<T, N, HASH, EQUAL>'''
+
+    def try_read_std_optional(self, slot, entry):
         try:
             # libstdc++
-            v = value['_M_payload']['_M_payload']['_M_value']
+            v = entry['_M_payload']['_M_payload']['_M_value']
             return slot, v
-            # return str(kvp['key']), kvp['value']
         except:
             return None
 
@@ -213,33 +219,16 @@
     'UtilsHashset', '^tint::utils::Hashset<.*>$', UtilsHashsetPrinter)
 
 
-class UtilsHashmapPrinter(Printer):
+class UtilsHashmapPrinter(UtilsHashmapBasePrinter):
     '''Printer for Hashmap<K, V, N, HASH, EQUAL>'''
 
-    def __init__(self, val):
-        super(UtilsHashmapPrinter, self).__init__(val)
-        self.hash_set = UtilsHashsetPrinter(self.val['set_'])
-        # Replace the lookup function so we can extract the key and value out of the std::optionals in the Hashset
-        self.hash_set.try_read_std_optional_func = self.try_read_std_optional
-
-    def to_string(self):
-        return self.hash_set.to_string()
-
-    def children(self):
-        return self.hash_set.children()
-
-    def display_hint(self):
-        return 'array'
-
-    def try_read_std_optional(self, slot, value):
+    def try_read_std_optional(self, slot, entry):
         try:
             # libstdc++
-            kvp = value['_M_payload']['_M_payload']['_M_value']
+            kvp = entry['_M_payload']['_M_payload']['_M_value']
             return str(kvp['key']), kvp['value']
         except:
-            pass
-        # Failed, fall back on hash_set
-        return self.hash_set.try_read_std_optional(slot, value)
+            return None
 
 
 pp_set.add_printer(
diff --git a/src/tint/tint_lldb.py b/src/tint/tint_lldb.py
index 85a93ff..84e6ad3 100644
--- a/src/tint/tint_lldb.py
+++ b/src/tint/tint_lldb.py
@@ -297,8 +297,8 @@
         return self.slice_printer.get_child_at_index(index)
 
 
-class UtilsHashsetPrinter(Printer):
-    '''Printer for Hashset<T, N, HASH, EQUAL>'''
+class UtilsHashmapBasePrinter(Printer):
+    '''Base Printer for HashmapBase-derived types'''
 
     def initialize(self):
         self.slice = UtilsVectorPrinter(
@@ -326,26 +326,33 @@
     def get_child_at_index(self, index):
         slot = self.valid_slots[index]
         v = self.slice.value_at(slot)
-        value = member(v, 'value')
+        entry = member(v, 'entry')
 
-        # value is a std::optional, let's try to extract its value for display
-        kvp = self.try_read_std_optional_func(slot, value)
+        # entry is a std::optional, let's try to extract its value for display
+        kvp = self.try_read_std_optional_func(slot, entry)
         if kvp is None:
-            # If we failed, just output the slot and value as is, which will use
+            # If we failed, just output the slot and entry as is, which will use
             # the default printer for std::optional.
-            kvp = slot, value
+            kvp = slot, entry
 
         return kvp[1].CreateChildAtOffset('[{}]'.format(kvp[0]), 0, kvp[1].GetType())
 
-    def try_read_std_optional(self, slot, value):
+    def try_read_std_optional(self, slot, entry):
+        return None
+
+
+class UtilsHashsetPrinter(UtilsHashmapBasePrinter):
+    '''Printer for Hashset<T, N, HASH, EQUAL>'''
+
+    def try_read_std_optional(self, slot, entry):
         try:
             # libc++
-            v = value.EvaluateExpression('__val_')
+            v = entry.EvaluateExpression('__val_')
             if v.name is not None:
                 return slot, v
 
             # libstdc++
-            v = value.EvaluateExpression('_M_payload._M_payload._M_value')
+            v = entry.EvaluateExpression('_M_payload._M_payload._M_value')
             if v.name is not None:
                 return slot, v
             return None
@@ -353,45 +360,24 @@
             return None
 
 
-class UtilsHashmapPrinter(Printer):
+class UtilsHashmapPrinter(UtilsHashsetPrinter):
     '''Printer for Hashmap<K, V, N, HASH, EQUAL>'''
 
-    def initialize(self):
-        self.hash_set = UtilsHashsetPrinter(self.member('set_'))
-        # Replace the lookup function so we can extract the key and value out of the std::optionals in the Hashset
-        self.hash_set.try_read_std_optional_func = self.try_read_std_optional
-
-    def update(self):
-        self.hash_set.update()
-
-    def get_summary(self):
-        return self.hash_set.get_summary()
-
-    def num_children(self):
-        return self.hash_set.num_children()
-
-    def has_children(self):
-        return self.hash_set.has_children()
-
-    def get_child_at_index(self, index):
-        return self.hash_set.get_child_at_index(index)
-
-    def try_read_std_optional(self, slot, value):
+    def try_read_std_optional(self, slot, entry):
         try:
             # libc++
-            val = value.EvaluateExpression('__val_')
+            val = entry.EvaluateExpression('__val_')
             k = val.EvaluateExpression('key')
             v = val.EvaluateExpression('value')
             if k.name is not None and v.name is not None:
                 return k.GetValue(), v
 
             # libstdc++
-            val = value.EvaluateExpression('_M_payload._M_payload._M_value')
+            val = entry.EvaluateExpression('_M_payload._M_payload._M_value')
             k = val.EvaluateExpression('key')
             v = val.EvaluateExpression('value')
             if k.name is not None and v.name is not None:
                 return k.GetValue(), v
+            return None
         except:
-            pass
-        # Failed, fall back on hash_set
-        return self.hash_set.try_read_std_optional(slot, value)
+            return None
diff --git a/src/tint/transform/builtin_polyfill.cc b/src/tint/transform/builtin_polyfill.cc
index e80436d..17fcc20 100644
--- a/src/tint/transform/builtin_polyfill.cc
+++ b/src/tint/transform/builtin_polyfill.cc
@@ -462,6 +462,14 @@
         auto name = b.Symbols().New("tint_insert_bits");
         uint32_t width = WidthOf(ty);
 
+        // Currently in WGSL parameters of insertBits must be i32, u32, vecN<i32> or vecN<u32>
+        if (!sem::Type::DeepestElementOf(ty)->IsAnyOf<sem::I32, sem::U32>()) {
+            TINT_ICE(Transform, b.Diagnostics())
+                << "insertBits polyfill only support i32, u32, and vector of i32 or u32, got "
+                << b.FriendlyName(ty);
+            return {};
+        }
+
         constexpr uint32_t W = 32u;  // 32-bit
 
         auto V = [&](auto value) -> const ast::Expression* {
@@ -481,21 +489,60 @@
             return b.vec(b.ty.u32(), width, value);
         };
 
-        utils::Vector<const ast::Statement*, 8> body = {
-            b.Decl(b.Let("s", b.Call("min", "offset", u32(W)))),
-            b.Decl(b.Let("e", b.Call("min", u32(W), b.Add("s", "count")))),
-        };
+        // Polyfill algorithm:
+        //      s = min(offset, 32u);
+        //      e = min(32u, (s + count));
+        //      mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+        //      return (((n << s) & mask) | (v & ~(mask)));
+        // Note that the algorithm above use the left-shifting in C++ manner, but in WGSL, HLSL, MSL
+        // the rhs are modulo to bit-width of lhs (that is 32u in this case), and in GLSL the result
+        // is undefined if rhs is greater than or equal to bit-width of lhs. The results of `x << y`
+        // in C++ and HLSL are different when `y >= 32u`, and the `s` and `e` defined above can be
+        // 32u, which are cases we must handle specially. Replace all `(x << y)` to
+        // `select(Tx(), x << y, y < 32u)`, in which `Tx` is the type of x, where y can be greater
+        // than or equal to 32u.
+        // WGSL polyfill function:
+        //      fn tint_insert_bits(v : T, n : T, offset : u32, count : u32) -> T {
+        //          let e = offset + count;
+        //          let mask = (
+        //                        (select(0u, 1u << offset, offset < 32u) - 1u) ^
+        //                        (select(0u, 1u << e, e < 32u) - 1u)
+        //                     );
+        //          return ((select(T(), n << offset, offset < 32u) & mask) | (v & ~(mask)));
+        //      }
+
+        utils::Vector<const ast::Statement*, 8> body;
 
         switch (polyfill.insert_bits) {
             case Level::kFull:
-                // let mask = ((1 << s) - 1) ^ ((1 << e) - 1)
+                // let e = offset + count;
+                body.Push(b.Decl(b.Let("e", b.Add("offset", "count"))));
+
+                // let mask = (
+                //              (select(0u, 1u << offset, offset < 32u) - 1u) ^
+                //              (select(0u, 1u << e, e < 32u) - 1u)
+                //            );
                 body.Push(b.Decl(b.Let(
-                    "mask", b.Xor(b.Sub(b.Shl(1_u, "s"), 1_u), b.Sub(b.Shl(1_u, "e"), 1_u)))));
-                // return ((n << s) & mask) | (v & ~mask)
-                body.Push(b.Return(b.Or(b.And(b.Shl("n", U("s")), V("mask")),
-                                        b.And("v", V(b.Complement("mask"))))));
+                    "mask",
+                    b.Xor(  //
+                        b.Sub(
+                            b.Call("select", 0_u, b.Shl(1_u, "offset"), b.LessThan("offset", 32_u)),
+                            1_u),
+                        b.Sub(b.Call("select", 0_u, b.Shl(1_u, "e"), b.LessThan("e", 32_u)),
+                              1_u)  //
+                        ))));
+
+                // return ((select(T(), n << offset, offset < 32u) & mask) | (v & ~(mask)));
+                body.Push(
+                    b.Return(b.Or(b.And(b.Call("select", b.Construct(T(ty)),
+                                               b.Shl("n", U("offset")), b.LessThan("offset", 32_u)),
+                                        V("mask")),
+                                  b.And("v", V(b.Complement("mask"))))));
+
                 break;
             case Level::kClampParameters:
+                body.Push(b.Decl(b.Let("s", b.Call("min", "offset", u32(W)))));
+                body.Push(b.Decl(b.Let("e", b.Call("min", u32(W), b.Add("s", "count")))));
                 body.Push(b.Return(b.Call("insertBits", "v", "n", "s", b.Sub("e", "s"))));
                 break;
             default:
diff --git a/src/tint/transform/builtin_polyfill_test.cc b/src/tint/transform/builtin_polyfill_test.cc
index 3b7a42c..1e380d4 100644
--- a/src/tint/transform/builtin_polyfill_test.cc
+++ b/src/tint/transform/builtin_polyfill_test.cc
@@ -1722,10 +1722,9 @@
 
     auto* expect = R"(
 fn tint_insert_bits(v : i32, n : i32, offset : u32, count : u32) -> i32 {
-  let s = min(offset, 32u);
-  let e = min(32u, (s + count));
-  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
-  return (((n << s) & i32(mask)) | (v & i32(~(mask))));
+  let e = (offset + count);
+  let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+  return ((select(i32(), (n << offset), (offset < 32u)) & i32(mask)) | (v & i32(~(mask))));
 }
 
 fn f() {
@@ -1749,10 +1748,9 @@
 
     auto* expect = R"(
 fn tint_insert_bits(v : u32, n : u32, offset : u32, count : u32) -> u32 {
-  let s = min(offset, 32u);
-  let e = min(32u, (s + count));
-  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
-  return (((n << s) & mask) | (v & ~(mask)));
+  let e = (offset + count);
+  let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+  return ((select(u32(), (n << offset), (offset < 32u)) & mask) | (v & ~(mask)));
 }
 
 fn f() {
@@ -1776,10 +1774,9 @@
 
     auto* expect = R"(
 fn tint_insert_bits(v : vec3<i32>, n : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
-  let s = min(offset, 32u);
-  let e = min(32u, (s + count));
-  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
-  return (((n << vec3<u32>(s)) & vec3<i32>(i32(mask))) | (v & vec3<i32>(i32(~(mask)))));
+  let e = (offset + count);
+  let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+  return ((select(vec3<i32>(), (n << vec3<u32>(offset)), (offset < 32u)) & vec3<i32>(i32(mask))) | (v & vec3<i32>(i32(~(mask)))));
 }
 
 fn f() {
@@ -1803,10 +1800,9 @@
 
     auto* expect = R"(
 fn tint_insert_bits(v : vec3<u32>, n : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
-  let s = min(offset, 32u);
-  let e = min(32u, (s + count));
-  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
-  return (((n << vec3<u32>(s)) & vec3<u32>(mask)) | (v & vec3<u32>(~(mask))));
+  let e = (offset + count);
+  let mask = ((select(0u, (1u << offset), (offset < 32u)) - 1u) ^ (select(0u, (1u << e), (e < 32u)) - 1u));
+  return ((select(vec3<u32>(), (n << vec3<u32>(offset)), (offset < 32u)) & vec3<u32>(mask)) | (v & vec3<u32>(~(mask))));
 }
 
 fn f() {
diff --git a/src/tint/transform/clamp_frag_depth_test.cc b/src/tint/transform/clamp_frag_depth_test.cc
index b94d5af..7c7019a 100644
--- a/src/tint/transform/clamp_frag_depth_test.cc
+++ b/src/tint/transform/clamp_frag_depth_test.cc
@@ -169,7 +169,7 @@
         @fragment fn main() -> @builtin(frag_depth) f32 {
             return 0.0;
         }
-        @fragment fn friend() -> @location(0) f32 {
+        @fragment fn other() -> @location(0) f32 {
             return 0.0;
         }
     )";
@@ -194,7 +194,7 @@
 }
 
 @fragment
-fn friend() -> @location(0) f32 {
+fn other() -> @location(0) f32 {
   return 0.0;
 }
 )";
diff --git a/src/tint/transform/multiplanar_external_texture_test.cc b/src/tint/transform/multiplanar_external_texture_test.cc
index 4416d35..8f2b014 100644
--- a/src/tint/transform/multiplanar_external_texture_test.cc
+++ b/src/tint/transform/multiplanar_external_texture_test.cc
@@ -580,9 +580,9 @@
 
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
-  var signed = textureLoad(ext_tex, vec2<i32>(1));
-  var unsigned = textureLoad(ext_tex, vec2<u32>(1));
-  return signed + unsigned;
+  var val_signed = textureLoad(ext_tex, vec2<i32>(1));
+  var val_unsigned = textureLoad(ext_tex, vec2<u32>(1));
+  return val_signed + val_unsigned;
 }
 )";
 
@@ -652,9 +652,9 @@
 
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
-  var signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
-  var unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
-  return (signed + unsigned);
+  var val_signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
+  var val_unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
+  return (val_signed + val_unsigned);
 }
 )";
 
@@ -670,9 +670,9 @@
     auto* src = R"(
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
-  var signed = textureLoad(ext_tex, vec2<i32>(1));
-  var unsigned = textureLoad(ext_tex, vec2<u32>(1));
-  return signed + unsigned;
+  var val_signed = textureLoad(ext_tex, vec2<i32>(1));
+  var val_unsigned = textureLoad(ext_tex, vec2<u32>(1));
+  return val_signed + val_unsigned;
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -742,9 +742,9 @@
 
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
-  var signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
-  var unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
-  return (signed + unsigned);
+  var val_signed = textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1), ext_tex_params);
+  var val_unsigned = textureLoadExternal_1(ext_tex, ext_tex_plane_1, vec2<u32>(1), ext_tex_params);
+  return (val_signed + val_unsigned);
 }
 
 @group(0) @binding(0) var ext_tex : texture_2d<f32>;
diff --git a/src/tint/transform/robustness_test.cc b/src/tint/transform/robustness_test.cc
index 990bbde..8b9a8ad 100644
--- a/src/tint/transform/robustness_test.cc
+++ b/src/tint/transform/robustness_test.cc
@@ -787,7 +787,7 @@
 @group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
 @group(0) @binding(0) var tex_external : texture_external;
 
-fn signed() {
+fn idx_signed() {
   var array_idx : i32;
   var level_idx : i32;
   var sample_idx : i32;
@@ -802,7 +802,7 @@
   textureLoad(tex_external, vec2<i32>(1, 2));
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   var array_idx : u32;
   var level_idx : u32;
   var sample_idx : u32;
@@ -836,7 +836,7 @@
 
 @group(0) @binding(0) var tex_external : texture_external;
 
-fn signed() {
+fn idx_signed() {
   var array_idx : i32;
   var level_idx : i32;
   var sample_idx : i32;
@@ -850,7 +850,7 @@
   textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   var array_idx : u32;
   var level_idx : u32;
   var sample_idx : u32;
@@ -873,7 +873,7 @@
 // Clamp textureLoad() coord, array_index and level values
 TEST_F(RobustnessTest, TextureLoad_Clamp_OutOfOrder) {
     auto* src = R"(
-fn signed() {
+fn idx_signed() {
   var array_idx : i32;
   var level_idx : i32;
   var sample_idx : i32;
@@ -888,7 +888,7 @@
   textureLoad(tex_external, vec2<i32>(1, 2));
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   var array_idx : u32;
   var level_idx : u32;
   var sample_idx : u32;
@@ -915,7 +915,7 @@
 
     auto* expect =
         R"(
-fn signed() {
+fn idx_signed() {
   var array_idx : i32;
   var level_idx : i32;
   var sample_idx : i32;
@@ -929,7 +929,7 @@
   textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   var array_idx : u32;
   var level_idx : u32;
   var sample_idx : u32;
@@ -976,14 +976,14 @@
 
 @group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
 
-fn signed() {
+fn idx_signed() {
   textureStore(tex1d, 10i, vec4<i32>());
   textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
   textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
   textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   textureStore(tex1d, 10u, vec4<i32>());
   textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
   textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
@@ -1000,14 +1000,14 @@
 
 @group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
 
-fn signed() {
+fn idx_signed() {
   textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
   textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
   textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
   textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
   textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
   textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
@@ -1023,14 +1023,14 @@
 // Clamp textureStore() coord, array_index and level values
 TEST_F(RobustnessTest, TextureStore_Clamp_OutOfOrder) {
     auto* src = R"(
-fn signed() {
+fn idx_signed() {
   textureStore(tex1d, 10i, vec4<i32>());
   textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
   textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
   textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   textureStore(tex1d, 10u, vec4<i32>());
   textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
   textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
@@ -1048,14 +1048,14 @@
 )";
 
     auto* expect = R"(
-fn signed() {
+fn idx_signed() {
   textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
   textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
   textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
   textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
 }
 
-fn unsigned() {
+fn idx_unsigned() {
   textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
   textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
   textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
diff --git a/src/tint/transform/single_entry_point_test.cc b/src/tint/transform/single_entry_point_test.cc
index 7451090..e9fa68f 100644
--- a/src/tint/transform/single_entry_point_test.cc
+++ b/src/tint/transform/single_entry_point_test.cc
@@ -240,32 +240,6 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(SingleEntryPointTest, WorkgroupSizeLetPreserved) {
-    auto* src = R"(
-let size : i32 = 1;
-
-@compute @workgroup_size(size)
-fn main() {
-}
-)";
-
-    auto* expect = R"(
-const size : i32 = 1;
-
-@compute @workgroup_size(size)
-fn main() {
-}
-)";
-
-    SingleEntryPoint::Config cfg("main");
-
-    DataMap data;
-    data.Add<SingleEntryPoint::Config>(cfg);
-    auto got = Run<SingleEntryPoint>(src, data);
-
-    EXPECT_EQ(expect, str(got));
-}
-
 TEST_F(SingleEntryPointTest, WorkgroupSizeConstPreserved) {
     auto* src = R"(
 const size : i32 = 1;
diff --git a/src/tint/transform/unshadow_test.cc b/src/tint/transform/unshadow_test.cc
index f5a8102..c731e76 100644
--- a/src/tint/transform/unshadow_test.cc
+++ b/src/tint/transform/unshadow_test.cc
@@ -384,82 +384,6 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(UnshadowTest, LocalShadowsGlobalLet) {
-    auto* src = R"(
-let a : i32 = 1;
-
-fn X() {
-  var a = (a == 123);
-}
-
-fn Y() {
-  let a = (a == 321);
-}
-
-fn Z() {
-  const a = 321;
-}
-)";
-
-    auto* expect = R"(
-const a : i32 = 1;
-
-fn X() {
-  var a_1 = (a == 123);
-}
-
-fn Y() {
-  let a_2 = (a == 321);
-}
-
-fn Z() {
-  const a_3 = 321;
-}
-)";
-
-    auto got = Run<Unshadow>(src);
-
-    EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(UnshadowTest, LocalShadowsGlobalLet_OutOfOrder) {
-    auto* src = R"(
-fn X() {
-  var a = (a == 123);
-}
-
-fn Y() {
-  let a = (a == 321);
-}
-
-fn Z() {
-  const a = 321;
-}
-
-let a : i32 = 1;
-)";
-
-    auto* expect = R"(
-fn X() {
-  var a_1 = (a == 123);
-}
-
-fn Y() {
-  let a_2 = (a == 321);
-}
-
-fn Z() {
-  const a_3 = 321;
-}
-
-const a : i32 = 1;
-)";
-
-    auto got = Run<Unshadow>(src);
-
-    EXPECT_EQ(expect, str(got));
-}
-
 TEST_F(UnshadowTest, LocalShadowsGlobalConst) {
     auto* src = R"(
 const a : i32 = 1;
@@ -732,46 +656,6 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(UnshadowTest, ParamShadowsGlobalLet) {
-    auto* src = R"(
-let a : i32 = 1;
-
-fn F(a : bool) {
-}
-)";
-
-    auto* expect = R"(
-const a : i32 = 1;
-
-fn F(a_1 : bool) {
-}
-)";
-
-    auto got = Run<Unshadow>(src);
-
-    EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(UnshadowTest, ParamShadowsGlobalLet_OutOfOrder) {
-    auto* src = R"(
-fn F(a : bool) {
-}
-
-let a : i32 = 1;
-)";
-
-    auto* expect = R"(
-fn F(a_1 : bool) {
-}
-
-const a : i32 = 1;
-)";
-
-    auto got = Run<Unshadow>(src);
-
-    EXPECT_EQ(expect, str(got));
-}
-
 TEST_F(UnshadowTest, ParamShadowsGlobalConst) {
     auto* src = R"(
 const a : i32 = 1;