diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index c2380c36..4ea7ab7 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -108,6 +108,8 @@
             return "vector_times_matrix";
         case BuiltinFn::kVectorTimesScalar:
             return "vector_times_scalar";
+        case BuiltinFn::kNormalize:
+            return "normalize";
         case BuiltinFn::kSdot:
             return "sdot";
         case BuiltinFn::kUdot:
@@ -161,6 +163,7 @@
         case BuiltinFn::kSdot:
         case BuiltinFn::kUdot:
         case BuiltinFn::kNone:
+        case BuiltinFn::kNormalize:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index a539fa5..cba1b82 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -72,6 +72,7 @@
         case BuiltinFn::kSdot:
         case BuiltinFn::kUdot:
         case BuiltinFn::kNone:
+        case BuiltinFn::kNormalize:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 8e582ab..0e49a9b 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -81,6 +81,7 @@
     kSelect,
     kVectorTimesMatrix,
     kVectorTimesScalar,
+    kNormalize,
     kSdot,
     kUdot,
     kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 5b40302..8e9b6d0 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -4881,6 +4881,28 @@
   {
     /* [149] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(21),
+    /* parameters */ ParameterIndex(4),
+    /* return_matcher_indices */ MatcherIndicesIndex(3),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [150] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 2,
+    /* templates */ TemplateIndex(42),
+    /* parameters */ ParameterIndex(279),
+    /* return_matcher_indices */ MatcherIndicesIndex(25),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [151] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
@@ -4890,7 +4912,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [150] */
+    /* [152] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 4,
     /* num_explicit_templates */ 0,
@@ -4901,7 +4923,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [151] */
+    /* [153] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 6,
     /* num_explicit_templates */ 0,
@@ -4912,7 +4934,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [152] */
+    /* [154] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -4923,7 +4945,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [153] */
+    /* [155] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 4,
     /* num_explicit_templates */ 0,
@@ -4934,7 +4956,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [154] */
+    /* [156] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4945,7 +4967,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [155] */
+    /* [157] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4956,7 +4978,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [156] */
+    /* [158] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4967,7 +4989,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [157] */
+    /* [159] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4978,7 +5000,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [158] */
+    /* [160] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4989,7 +5011,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [159] */
+    /* [161] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -5000,7 +5022,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [160] */
+    /* [162] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -5011,7 +5033,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [161] */
+    /* [163] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -5031,91 +5053,91 @@
     /* [0] */
     /* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(149),
+    /* overloads */ OverloadIndex(151),
   },
   {
     /* [1] */
     /* fn atomic_and[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [2] */
     /* fn atomic_compare_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, U, T, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(151),
+    /* overloads */ OverloadIndex(153),
   },
   {
     /* [3] */
     /* fn atomic_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [4] */
     /* fn atomic_iadd[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [5] */
     /* fn atomic_isub[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [6] */
     /* fn atomic_load[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(152),
+    /* overloads */ OverloadIndex(154),
   },
   {
     /* [7] */
     /* fn atomic_or[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [8] */
     /* fn atomic_smax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [9] */
     /* fn atomic_smin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [10] */
     /* fn atomic_store[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(153),
+    /* overloads */ OverloadIndex(155),
   },
   {
     /* [11] */
     /* fn atomic_umax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [12] */
     /* fn atomic_umin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [13] */
     /* fn atomic_xor[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(150),
+    /* overloads */ OverloadIndex(152),
   },
   {
     /* [14] */
     /* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(154),
+    /* overloads */ OverloadIndex(156),
   },
   {
     /* [15] */
@@ -5313,19 +5335,19 @@
     /* [26] */
     /* fn matrix_times_matrix[T : f32_f16, K : num, C : num, R : num](mat<K, R, T>, mat<C, K, T>) -> mat<C, R, T> */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(155),
+    /* overloads */ OverloadIndex(157),
   },
   {
     /* [27] */
     /* fn matrix_times_scalar[T : f32_f16, N : num, M : num](mat<N, M, T>, T) -> mat<N, M, T> */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(156),
+    /* overloads */ OverloadIndex(158),
   },
   {
     /* [28] */
     /* fn matrix_times_vector[T : f32_f16, N : num, M : num](mat<N, M, T>, vec<N, T>) -> vec<M, T> */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [29] */
@@ -5353,25 +5375,32 @@
     /* [31] */
     /* fn vector_times_matrix[T : f32_f16, N : num, M : num](vec<N, T>, mat<M, N, T>) -> vec<M, T> */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(158),
+    /* overloads */ OverloadIndex(160),
   },
   {
     /* [32] */
     /* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(159),
+    /* overloads */ OverloadIndex(161),
   },
   {
     /* [33] */
-    /* fn sdot(u32, u32, u32) -> i32 */
-    /* num overloads */ 1,
-    /* overloads */ OverloadIndex(160),
+    /* fn normalize[T : f32_f16](T) -> T */
+    /* fn normalize[N : num, T : f32_f16](vec<N, T>) -> vec<N, T> */
+    /* num overloads */ 2,
+    /* overloads */ OverloadIndex(149),
   },
   {
     /* [34] */
+    /* fn sdot(u32, u32, u32) -> i32 */
+    /* num overloads */ 1,
+    /* overloads */ OverloadIndex(162),
+  },
+  {
+    /* [35] */
     /* fn udot(u32, u32, u32) -> u32 */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(161),
+    /* overloads */ OverloadIndex(163),
   },
 };
 
diff --git a/src/tint/lang/spirv/reader/helper_test.h b/src/tint/lang/spirv/reader/helper_test.h
index 7bb17bf..0d14aca 100644
--- a/src/tint/lang/spirv/reader/helper_test.h
+++ b/src/tint/lang/spirv/reader/helper_test.h
@@ -28,14 +28,12 @@
 #ifndef SRC_TINT_LANG_SPIRV_READER_HELPER_TEST_H_
 #define SRC_TINT_LANG_SPIRV_READER_HELPER_TEST_H_
 
-#include <iostream>
 #include <string>
-#include <vector>
 
 #include "gmock/gmock.h"
 #include "gtest/gtest.h"
+
 #include "src/tint/lang/core/ir/disassembler.h"
-#include "src/tint/lang/core/ir/module.h"
 #include "src/tint/lang/core/ir/validator.h"
 #include "src/tint/lang/spirv/reader/common/helper_test.h"
 #include "src/tint/lang/spirv/reader/reader.h"
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
index 0434bc1..47b240a 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -832,7 +832,7 @@
 // So we have to test it separately, as it does not fit the patterns tested
 // above.
 
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Scalar) {
+TEST_F(SpirvReaderTest, Normalize_Scalar) {
     // Scalar normalize maps to sign.
     EXPECT_IR(Preamble() + R"(
      %1 = OpExtInst %float %glsl Normalize %float_50
@@ -851,7 +851,7 @@
 )");
 }
 
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector2) {
+TEST_F(SpirvReaderTest, Normalize_Vector2) {
     EXPECT_IR(Preamble() + R"(
      %1 = OpExtInst %v2float %glsl Normalize %v2float_50_60
      %2 = OpCopyObject %v2float %1
@@ -869,7 +869,7 @@
 )");
 }
 
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector3) {
+TEST_F(SpirvReaderTest, Normalize_Vector3) {
     EXPECT_IR(Preamble() + R"(
      %1 = OpExtInst %v3float %glsl Normalize %v3float_50_60_70
      %2 = OpCopyObject %v3float %1
@@ -887,7 +887,7 @@
 )");
 }
 
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector4) {
+TEST_F(SpirvReaderTest, Normalize_Vector4) {
     EXPECT_IR(Preamble() + R"(
      %1 = OpExtInst %v4float %glsl Normalize %v4float_50_50_50_50
      %2 = OpCopyObject %v4float %1
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
index b29d625..eac8130 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -39,11 +39,13 @@
 cc_library(
   name = "lower",
   srcs = [
+    "builtins.cc",
     "lower.cc",
     "shader_io.cc",
     "vector_element_pointer.cc",
   ],
   hdrs = [
+    "builtins.h",
     "lower.h",
     "shader_io.h",
     "vector_element_pointer.h",
@@ -55,6 +57,9 @@
     "//src/tint/lang/core/intrinsic",
     "//src/tint/lang/core/ir",
     "//src/tint/lang/core/type",
+    "//src/tint/lang/spirv",
+    "//src/tint/lang/spirv/intrinsic",
+    "//src/tint/lang/spirv/ir",
     "//src/tint/utils",
     "//src/tint/utils/containers",
     "//src/tint/utils/diagnostic",
@@ -75,6 +80,7 @@
   name = "test",
   alwayslink = True,
   srcs = [
+    "builtins_test.cc",
     "shader_io_test.cc",
     "vector_element_pointer_test.cc",
   ],
@@ -86,6 +92,9 @@
     "//src/tint/lang/core/ir",
     "//src/tint/lang/core/ir/transform:test",
     "//src/tint/lang/core/type",
+    "//src/tint/lang/spirv",
+    "//src/tint/lang/spirv/intrinsic",
+    "//src/tint/lang/spirv/ir",
     "//src/tint/lang/spirv/reader/lower",
     "//src/tint/utils",
     "//src/tint/utils/containers",
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
index 6869be8..e39b8bb 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -39,6 +39,8 @@
 # Kind:      lib
 ################################################################################
 tint_add_target(tint_lang_spirv_reader_lower lib
+  lang/spirv/reader/lower/builtins.cc
+  lang/spirv/reader/lower/builtins.h
   lang/spirv/reader/lower/lower.cc
   lang/spirv/reader/lower/lower.h
   lang/spirv/reader/lower/shader_io.cc
@@ -54,6 +56,9 @@
   tint_lang_core_intrinsic
   tint_lang_core_ir
   tint_lang_core_type
+  tint_lang_spirv
+  tint_lang_spirv_intrinsic
+  tint_lang_spirv_ir
   tint_utils
   tint_utils_containers
   tint_utils_diagnostic
@@ -76,6 +81,7 @@
 # Kind:      test
 ################################################################################
 tint_add_target(tint_lang_spirv_reader_lower_test test
+  lang/spirv/reader/lower/builtins_test.cc
   lang/spirv/reader/lower/shader_io_test.cc
   lang/spirv/reader/lower/vector_element_pointer_test.cc
 )
@@ -88,6 +94,9 @@
   tint_lang_core_ir
   tint_lang_core_ir_transform_test
   tint_lang_core_type
+  tint_lang_spirv
+  tint_lang_spirv_intrinsic
+  tint_lang_spirv_ir
   tint_lang_spirv_reader_lower
   tint_utils
   tint_utils_containers
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
index 6012672..37bc08d 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -45,6 +45,8 @@
 
 libtint_source_set("lower") {
   sources = [
+    "builtins.cc",
+    "builtins.h",
     "lower.cc",
     "lower.h",
     "shader_io.cc",
@@ -60,6 +62,9 @@
     "${tint_src_dir}/lang/core/intrinsic",
     "${tint_src_dir}/lang/core/ir",
     "${tint_src_dir}/lang/core/type",
+    "${tint_src_dir}/lang/spirv",
+    "${tint_src_dir}/lang/spirv/intrinsic",
+    "${tint_src_dir}/lang/spirv/ir",
     "${tint_src_dir}/utils",
     "${tint_src_dir}/utils/containers",
     "${tint_src_dir}/utils/diagnostic",
@@ -76,6 +81,7 @@
 if (tint_build_unittests) {
   tint_unittests_source_set("unittests") {
     sources = [
+      "builtins_test.cc",
       "shader_io_test.cc",
       "vector_element_pointer_test.cc",
     ]
@@ -89,6 +95,9 @@
       "${tint_src_dir}/lang/core/ir",
       "${tint_src_dir}/lang/core/ir/transform:unittests",
       "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/spirv",
+      "${tint_src_dir}/lang/spirv/intrinsic",
+      "${tint_src_dir}/lang/spirv/ir",
       "${tint_src_dir}/lang/spirv/reader/lower",
       "${tint_src_dir}/utils",
       "${tint_src_dir}/utils/containers",
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
new file mode 100644
index 0000000..c9edc17
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -0,0 +1,100 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
+
+namespace tint::spirv::reader::lower {
+
+namespace {
+
+using namespace tint::core::fluent_types;  // NOLINT
+
+/// PIMPL state for the transform.
+struct State {
+    /// The IR module.
+    core::ir::Module& ir;
+
+    /// The IR builder.
+    core::ir::Builder b{ir};
+
+    /// The type manager.
+    core::type::Manager& ty{ir.Types()};
+
+    /// Process the module.
+    void Process() {
+        Vector<spirv::ir::BuiltinCall*, 4> builtin_worklist;
+        for (auto* inst : ir.Instructions()) {
+            if (auto* builtin = inst->As<spirv::ir::BuiltinCall>()) {
+                builtin_worklist.Push(builtin);
+            }
+        }
+
+        // Replace the builtins that we found.
+        for (auto* builtin : builtin_worklist) {
+            switch (builtin->Func()) {
+                case spirv::BuiltinFn::kNormalize:
+                    Normalize(builtin);
+                    break;
+                default:
+                    TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
+            }
+        }
+    }
+
+    void Normalize(spirv::ir::BuiltinCall* call) {
+        auto* arg = call->Args()[0];
+
+        b.InsertBefore(call, [&] {
+            core::BuiltinFn fn = core::BuiltinFn::kNormalize;
+            if (arg->Type()->IsScalar()) {
+                fn = core::BuiltinFn::kSign;
+            }
+            b.CallWithResult(call->DetachResult(), fn, Vector<core::ir::Value*, 1>{arg});
+        });
+        call->Destroy();
+    }
+};
+
+}  // namespace
+
+Result<SuccessType> Builtins(core::ir::Module& ir) {
+    auto result = ValidateAndDumpIfNeeded(ir, "spirv.Builtins");
+    if (result != Success) {
+        return result.Failure();
+    }
+
+    State{ir}.Process();
+
+    return Success;
+}
+
+}  // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/lower/builtins.h b/src/tint/lang/spirv/reader/lower/builtins.h
new file mode 100644
index 0000000..f3b8009
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins.h
@@ -0,0 +1,48 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
+#define SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::reader::lower {
+
+/// Builtins is a transform that converts SPIR-V builtin methods into the form expected by Tint's
+/// core IR.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> Builtins(core::ir::Module& module);
+
+}  // namespace tint::spirv::reader::lower
+
+#endif  // SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
new file mode 100644
index 0000000..95c814e
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -0,0 +1,103 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
+
+namespace tint::spirv::reader::lower {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using SpirvParser_BuiltinsTest = core::ir::transform::TransformTest;
+
+TEST_F(SpirvParser_BuiltinsTest, Normalize_Scalar) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.Call<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kNormalize, 10_f);
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = spirv.normalize 10.0f
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = sign 10.0f
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, Normalize_Vector) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.Call<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kNormalize,
+                                       b.Splat(ty.vec2<f32>(), 10_f));
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = spirv.normalize vec2<f32>(10.0f)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = normalize vec2<f32>(10.0f)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
+}  // namespace
+}  // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/lower/lower.cc b/src/tint/lang/spirv/reader/lower/lower.cc
index f1861d4..2eecf38 100644
--- a/src/tint/lang/spirv/reader/lower/lower.cc
+++ b/src/tint/lang/spirv/reader/lower/lower.cc
@@ -28,6 +28,7 @@
 #include "src/tint/lang/spirv/reader/lower/lower.h"
 
 #include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
 #include "src/tint/lang/spirv/reader/lower/shader_io.h"
 #include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
 
@@ -44,6 +45,7 @@
 
     RUN_TRANSFORM(lower::VectorElementPointer, mod);
     RUN_TRANSFORM(lower::ShaderIO, mod);
+    RUN_TRANSFORM(lower::Builtins, mod);
 
     if (auto res = core::ir::ValidateAndDumpIfNeeded(mod, "spirv.Lower"); res != Success) {
         return res.Failure();
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 1fcad15..e99e6d0 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -51,6 +51,9 @@
     "//src/tint/lang/core/intrinsic",
     "//src/tint/lang/core/ir",
     "//src/tint/lang/core/type",
+    "//src/tint/lang/spirv",
+    "//src/tint/lang/spirv/intrinsic",
+    "//src/tint/lang/spirv/ir",
     "//src/tint/utils",
     "//src/tint/utils/containers",
     "//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index e8ddda8..d4c22cb 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -52,6 +52,9 @@
   tint_lang_core_intrinsic
   tint_lang_core_ir
   tint_lang_core_type
+  tint_lang_spirv
+  tint_lang_spirv_intrinsic
+  tint_lang_spirv_ir
   tint_utils
   tint_utils_containers
   tint_utils_diagnostic
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index c54f228..9f4c271 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -56,6 +56,9 @@
       "${tint_src_dir}/lang/core/intrinsic",
       "${tint_src_dir}/lang/core/ir",
       "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/spirv",
+      "${tint_src_dir}/lang/spirv/intrinsic",
+      "${tint_src_dir}/lang/spirv/ir",
       "${tint_src_dir}/utils",
       "${tint_src_dir}/utils/containers",
       "${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 3a70f3b..a32a8fb 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -46,6 +46,8 @@
 
 #include "src/tint/lang/core/ir/builder.h"
 #include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/spirv/builtin_fn.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
 #include "src/tint/lang/spirv/validate/validate.h"
 
 using namespace tint::core::fluent_types;  // NOLINT
@@ -633,7 +635,7 @@
     // Returns the WGSL standard library function for the given GLSL.std.450 extended instruction
     // operation code. This handles GLSL functions which directly translate to the WGSL equivalent.
     // Any non-direct translation is returned as `kNone`.
-    core::BuiltinFn GetGlslStd450FuncName(uint32_t ext_opcode) {
+    core::BuiltinFn GetGlslStd450WgslEquivalentFuncName(uint32_t ext_opcode) {
         switch (ext_opcode) {
             case GLSLstd450Acos:
                 return core::BuiltinFn::kAcos;
@@ -753,22 +755,40 @@
         return core::BuiltinFn::kNone;
     }
 
+    spirv::BuiltinFn GetGlslStd450SpirvEquivalentFuncName(uint32_t ext_opcode) {
+        switch (ext_opcode) {
+            case GLSLstd450Normalize:
+                return spirv::BuiltinFn::kNormalize;
+            default:
+                break;
+        }
+        return spirv::BuiltinFn::kNone;
+    }
+
     /// @param inst the SPIR-V instruction for OpAccessChain
     void EmitGlslStd450ExtInst(const spvtools::opt::Instruction& inst) {
         const auto ext_opcode = inst.GetSingleWordInOperand(1);
         auto* result_ty = Type(inst.type_id());
 
-        const auto fn = GetGlslStd450FuncName(ext_opcode);
-        if (fn == core::BuiltinFn::kNone) {
-            TINT_UNIMPLEMENTED() << "unhandled GLSL.std.450 instruction " << ext_opcode;
-        }
-
         Vector<core::ir::Value*, 4> operands;
         // All parameters to GLSL.std.450 extended instructions are IDs.
         for (uint32_t idx = 2; idx < inst.NumInOperands(); ++idx) {
             operands.Push(Value(inst.GetSingleWordInOperand(idx)));
         }
-        Emit(b_.Call(result_ty, fn, operands), inst.result_id());
+
+        const auto wgsl_fn = GetGlslStd450WgslEquivalentFuncName(ext_opcode);
+        if (wgsl_fn != core::BuiltinFn::kNone) {
+            Emit(b_.Call(result_ty, wgsl_fn, operands), inst.result_id());
+            return;
+        }
+
+        const auto spv_fn = GetGlslStd450SpirvEquivalentFuncName(ext_opcode);
+        if (spv_fn != spirv::BuiltinFn::kNone) {
+            Emit(b_.Call<spirv::ir::BuiltinCall>(result_ty, spv_fn, operands), inst.result_id());
+            return;
+        }
+
+        TINT_UNIMPLEMENTED() << "unhandled GLSL.std.450 instruction " << ext_opcode;
     }
 
     /// @param inst the SPIR-V instruction for OpAccessChain
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 0208789..cf7b5fc 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -318,6 +318,9 @@
 implicit(T: f32_f16, N: num, M: num) fn vector_times_matrix(vec<N, T>, mat<M, N, T>) -> vec<M, T>
 implicit(T: f32_f16, N: num) fn vector_times_scalar(vec<N, T>, T) -> vec<N, T>
 
+implicit(T: f32_f16) fn normalize(T) -> T
+implicit(N: num, T: f32_f16) fn normalize(vec<N, T>) -> vec<N, T>
+
 ////////////////////////////////////////////////////////////////////////////////
 // SPV_KHR_integer_dot_product instructions
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 7dbdd6c..b69b23c 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -116,6 +116,7 @@
 namespace {
 
 constexpr uint32_t kWriterVersion = 1;
+constexpr const char* kGLSLstd450 = "GLSL.std.450";
 
 SpvStorageClass StorageClass(core::AddressSpace addrspace) {
     switch (addrspace) {
@@ -1271,6 +1272,11 @@
     void EmitSpirvBuiltinCall(spirv::ir::BuiltinCall* builtin) {
         auto id = Value(builtin);
 
+        OperandList operands;
+        if (!builtin->Result(0)->Type()->Is<core::type::Void>()) {
+            operands = {Type(builtin->Result(0)->Type()), id};
+        }
+
         spv::Op op = spv::Op::Max;
         switch (builtin->Func()) {
             case spirv::BuiltinFn::kArrayLength:
@@ -1362,6 +1368,11 @@
             case spirv::BuiltinFn::kMatrixTimesVector:
                 op = spv::Op::OpMatrixTimesVector;
                 break;
+            case spirv::BuiltinFn::kNormalize:
+                op = spv::Op::OpExtInst;
+                operands.push_back(ImportGlslStd450());
+                operands.push_back(U32Operand(GLSLstd450Normalize));
+                break;
             case spirv::BuiltinFn::kSampledImage:
                 op = spv::Op::OpSampledImage;
                 break;
@@ -1390,16 +1401,21 @@
                 TINT_ICE() << "undefined spirv ir function";
         }
 
-        OperandList operands;
-        if (!builtin->Result(0)->Type()->Is<core::type::Void>()) {
-            operands = {Type(builtin->Result(0)->Type()), id};
-        }
         for (auto* arg : builtin->Args()) {
             operands.push_back(Value(arg));
         }
         current_function_.PushInst(op, operands);
     }
 
+    uint32_t ImportGlslStd450() {
+        return imports_.GetOrAdd(kGLSLstd450, [&] {
+            // Import the instruction set the first time it is requested.
+            auto import = module_.NextId();
+            module_.PushExtImport(spv::Op::OpExtInstImport, {import, Operand(kGLSLstd450)});
+            return import;
+        });
+    }
+
     /// Emit a builtin function call instruction.
     /// @param builtin the builtin call instruction to emit
     void EmitCoreBuiltinCall(core::ir::CoreBuiltinCall* builtin) {
@@ -1426,14 +1442,8 @@
 
         // Helper to set up the opcode and operand list for a GLSL extended instruction.
         auto glsl_ext_inst = [&](enum GLSLstd450 inst) {
-            constexpr const char* kGLSLstd450 = "GLSL.std.450";
             op = spv::Op::OpExtInst;
-            operands.push_back(imports_.GetOrAdd(kGLSLstd450, [&] {
-                // Import the instruction set the first time it is requested.
-                auto import = module_.NextId();
-                module_.PushExtImport(spv::Op::OpExtInstImport, {import, Operand(kGLSLstd450)});
-                return import;
-            }));
+            operands.push_back(ImportGlslStd450());
             operands.push_back(U32Operand(inst));
         };