diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index 6f1e558..302b464 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -50,6 +50,7 @@
     "//src/tint/lang/core/ir:test",
     "//src/tint/lang/core/type:test",
     "//src/tint/lang/core:test",
+    "//src/tint/lang/msl/ir:test",
     "//src/tint/lang/spirv/ir:test",
     "//src/tint/lang/wgsl/ast:test",
     "//src/tint/lang/wgsl/features:test",
@@ -107,6 +108,7 @@
       "//src/tint/lang/msl/writer/ast_printer:test",
       "//src/tint/lang/msl/writer/common:test",
       "//src/tint/lang/msl/writer/printer:test",
+      "//src/tint/lang/msl/writer/raise:test",
     ],
     "//conditions:default": [],
   }) + select({
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 670c21f..a231ddd 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -51,6 +51,7 @@
   tint_lang_core_ir_test
   tint_lang_core_type_test
   tint_lang_core_test
+  tint_lang_msl_ir_test
   tint_lang_spirv_ir_test
   tint_lang_wgsl_ast_test
   tint_lang_wgsl_features_test
@@ -117,6 +118,7 @@
     tint_lang_msl_writer_ast_printer_test
     tint_lang_msl_writer_common_test
     tint_lang_msl_writer_printer_test
+    tint_lang_msl_writer_raise_test
   )
 endif(TINT_BUILD_MSL_WRITER)
 
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index 17410ae..4fcb365 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -56,6 +56,7 @@
       "${tint_src_dir}/lang/core/ir:unittests",
       "${tint_src_dir}/lang/core/ir/transform:unittests",
       "${tint_src_dir}/lang/core/type:unittests",
+      "${tint_src_dir}/lang/msl/ir:unittests",
       "${tint_src_dir}/lang/spirv/ir:unittests",
       "${tint_src_dir}/lang/wgsl:unittests",
       "${tint_src_dir}/lang/wgsl/ast:unittests",
@@ -114,6 +115,7 @@
         "${tint_src_dir}/lang/msl/writer/ast_printer:unittests",
         "${tint_src_dir}/lang/msl/writer/common:unittests",
         "${tint_src_dir}/lang/msl/writer/printer:unittests",
+        "${tint_src_dir}/lang/msl/writer/raise:unittests",
       ]
     }
 
diff --git a/src/tint/lang/msl/BUILD.bazel b/src/tint/lang/msl/BUILD.bazel
index 1a114c6..af698ff 100644
--- a/src/tint/lang/msl/BUILD.bazel
+++ b/src/tint/lang/msl/BUILD.bazel
@@ -36,4 +36,19 @@
 
 load("//src/tint:flags.bzl", "COPTS")
 load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+  name = "msl",
+  srcs = [
+    "builtin_fn.cc",
+  ],
+  hdrs = [
+    "barrier_type.h",
+    "builtin_fn.h",
+  ],
+  deps = [
+    "//src/tint/utils/traits",
+  ],
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
 
diff --git a/src/tint/lang/msl/BUILD.cmake b/src/tint/lang/msl/BUILD.cmake
index bbcbda1..40cbc8e 100644
--- a/src/tint/lang/msl/BUILD.cmake
+++ b/src/tint/lang/msl/BUILD.cmake
@@ -34,5 +34,21 @@
 #                       Do not modify this file directly
 ################################################################################
 
+include(lang/msl/intrinsic/BUILD.cmake)
+include(lang/msl/ir/BUILD.cmake)
 include(lang/msl/validate/BUILD.cmake)
 include(lang/msl/writer/BUILD.cmake)
+
+################################################################################
+# Target:    tint_lang_msl
+# Kind:      lib
+################################################################################
+tint_add_target(tint_lang_msl lib
+  lang/msl/barrier_type.h
+  lang/msl/builtin_fn.cc
+  lang/msl/builtin_fn.h
+)
+
+tint_target_add_dependencies(tint_lang_msl lib
+  tint_utils_traits
+)
diff --git a/src/tint/lang/msl/BUILD.gn b/src/tint/lang/msl/BUILD.gn
index 6caa255..a092431 100644
--- a/src/tint/lang/msl/BUILD.gn
+++ b/src/tint/lang/msl/BUILD.gn
@@ -37,3 +37,12 @@
 import("../../../../scripts/tint_overrides_with_defaults.gni")
 
 import("${tint_src_dir}/tint.gni")
+
+libtint_source_set("msl") {
+  sources = [
+    "barrier_type.h",
+    "builtin_fn.cc",
+    "builtin_fn.h",
+  ]
+  deps = [ "${tint_src_dir}/utils/traits" ]
+}
diff --git a/src/tint/lang/msl/barrier_type.h b/src/tint/lang/msl/barrier_type.h
new file mode 100644
index 0000000..c3c13b4
--- /dev/null
+++ b/src/tint/lang/msl/barrier_type.h
@@ -0,0 +1,49 @@
+// Copyright 2023 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_MSL_BARRIER_TYPE_H_
+#define SRC_TINT_LANG_MSL_BARRIER_TYPE_H_
+
+#include <cstdint>
+
+namespace tint::msl {
+
+/// A barrier type
+enum BarrierType : uint8_t {
+    /// No type set
+    kNone = 0x0,
+    /// Device type
+    kDevice = 0x1,
+    /// Texture type
+    kTexture = 0x2,
+    /// Thead group type
+    kThreadGroup = 0x4,
+};
+
+}  // namespace tint::msl
+
+#endif  // SRC_TINT_LANG_MSL_BARRIER_TYPE_H_
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
new file mode 100644
index 0000000..18d96de
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -0,0 +1,51 @@
+// Copyright 2023 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/msl/builtin_fn.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl {
+
+const char* str(BuiltinFn i) {
+    switch (i) {
+        case BuiltinFn::kNone:
+            return "<none>";
+        case BuiltinFn::kThreadgroupBarrier:
+            return "threadgroup_barrier";
+    }
+    return "<unknown>";
+}
+
+}  // namespace tint::msl
diff --git a/src/tint/lang/msl/builtin_fn.cc.tmpl b/src/tint/lang/msl/builtin_fn.cc.tmpl
new file mode 100644
index 0000000..139045f
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.cc.tmpl
@@ -0,0 +1,31 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate builtin_fn.cc
+
+To update the generated file, run:
+    ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl {
+
+const char* str(BuiltinFn i) {
+    switch (i) {
+        case BuiltinFn::kNone:
+            return "<none>";
+{{- range $I.Sem.Builtins  }}
+        case BuiltinFn::k{{PascalCase .Name}}:
+            return "{{.Name}}";
+{{- end  }}
+    }
+    return "<unknown>";
+}
+
+}  // namespace tint::msl
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
new file mode 100644
index 0000000..c652488
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -0,0 +1,66 @@
+// Copyright 2023 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/msl/builtin_fn.h.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#ifndef SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+#define SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+
+#include <cstdint>
+#include <string>
+
+#include "src/tint/utils/traits/traits.h"
+
+// \cond DO_NOT_DOCUMENT
+namespace tint::msl {
+
+/// Enumerator of all builtin functions
+enum class BuiltinFn : uint8_t {
+    kThreadgroupBarrier,
+    kNone,
+};
+
+/// @returns the name of the builtin function type.
+const char* str(BuiltinFn i);
+
+/// Emits the name of the builtin function type.
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& o, BuiltinFn i) {
+    return o << str(i);
+}
+
+}  // namespace tint::msl
+// \endcond
+
+#endif  // SRC_TINT_LANG_MSL_BUILTIN_FN_H_
diff --git a/src/tint/lang/msl/builtin_fn.h.tmpl b/src/tint/lang/msl/builtin_fn.h.tmpl
new file mode 100644
index 0000000..491524f
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.h.tmpl
@@ -0,0 +1,47 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate builtin_fn.h
+
+To update the generated file, run:
+    ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+
+#ifndef SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+#define SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+
+#include <cstdint>
+#include <string>
+
+#include "src/tint/utils/traits/traits.h"
+
+// \cond DO_NOT_DOCUMENT
+namespace tint::msl {
+
+/// Enumerator of all builtin functions
+enum class BuiltinFn : uint8_t {
+{{- range $I.Sem.Builtins }}
+    k{{PascalCase .Name}},
+{{- end }}
+    kNone,
+};
+
+/// @returns the name of the builtin function type.
+const char* str(BuiltinFn i);
+
+/// Emits the name of the builtin function type.
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& o, BuiltinFn i) {
+  return o << str(i);
+}
+
+}  // namespace tint::msl
+// \endcond
+
+#endif  // SRC_TINT_LANG_MSL_BUILTIN_FN_H_
diff --git a/src/tint/lang/msl/intrinsic/BUILD.bazel b/src/tint/lang/msl/intrinsic/BUILD.bazel
new file mode 100644
index 0000000..de60b66
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.bazel
@@ -0,0 +1,69 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.bazel.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+load("//src/tint:flags.bzl", "COPTS")
+load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+  name = "intrinsic",
+  srcs = [
+    "data.cc",
+  ],
+  hdrs = [
+    "dialect.h",
+  ],
+  deps = [
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/type",
+    "//src/tint/lang/msl",
+    "//src/tint/utils/containers",
+    "//src/tint/utils/diagnostic",
+    "//src/tint/utils/ice",
+    "//src/tint/utils/id",
+    "//src/tint/utils/macros",
+    "//src/tint/utils/math",
+    "//src/tint/utils/memory",
+    "//src/tint/utils/result",
+    "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
+    "//src/tint/utils/text",
+    "//src/tint/utils/traits",
+  ],
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/msl/intrinsic/BUILD.cmake b/src/tint/lang/msl/intrinsic/BUILD.cmake
new file mode 100644
index 0000000..60cca31
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.cmake
@@ -0,0 +1,64 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.cmake.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+################################################################################
+# Target:    tint_lang_msl_intrinsic
+# Kind:      lib
+################################################################################
+tint_add_target(tint_lang_msl_intrinsic lib
+  lang/msl/intrinsic/data.cc
+  lang/msl/intrinsic/dialect.h
+)
+
+tint_target_add_dependencies(tint_lang_msl_intrinsic lib
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_type
+  tint_lang_msl
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
diff --git a/src/tint/lang/msl/intrinsic/BUILD.gn b/src/tint/lang/msl/intrinsic/BUILD.gn
new file mode 100644
index 0000000..61d0771
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.gn
@@ -0,0 +1,65 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.gn.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+import("../../../../../scripts/tint_overrides_with_defaults.gni")
+
+import("${tint_src_dir}/tint.gni")
+
+libtint_source_set("intrinsic") {
+  sources = [
+    "data.cc",
+    "dialect.h",
+  ]
+  deps = [
+    "${tint_src_dir}/lang/core",
+    "${tint_src_dir}/lang/core/constant",
+    "${tint_src_dir}/lang/core/intrinsic",
+    "${tint_src_dir}/lang/core/type",
+    "${tint_src_dir}/lang/msl",
+    "${tint_src_dir}/utils/containers",
+    "${tint_src_dir}/utils/diagnostic",
+    "${tint_src_dir}/utils/ice",
+    "${tint_src_dir}/utils/id",
+    "${tint_src_dir}/utils/macros",
+    "${tint_src_dir}/utils/math",
+    "${tint_src_dir}/utils/memory",
+    "${tint_src_dir}/utils/result",
+    "${tint_src_dir}/utils/rtti",
+    "${tint_src_dir}/utils/symbol",
+    "${tint_src_dir}/utils/text",
+    "${tint_src_dir}/utils/traits",
+  ]
+}
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
new file mode 100644
index 0000000..484b078
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -0,0 +1,188 @@
+// Copyright 2023 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/msl/intrinsic/data.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include <limits>
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/type_matchers.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/text/string_stream.h"
+
+namespace tint::msl::intrinsic {
+
+using namespace tint::core::intrinsic;  // NOLINT(build/namespaces)
+
+namespace {
+
+using ConstEvalFunctionIndex = tint::core::intrinsic::ConstEvalFunctionIndex;
+using IntrinsicInfo = tint::core::intrinsic::IntrinsicInfo;
+using MatchState = tint::core::intrinsic::MatchState;
+using Number = tint::core::intrinsic::Number;
+using NumberMatcher = tint::core::intrinsic::NumberMatcher;
+using NumberMatcherIndex = tint::core::intrinsic::NumberMatcherIndex;
+using NumberMatcherIndicesIndex = tint::core::intrinsic::NumberMatcherIndicesIndex;
+using OverloadFlag = tint::core::intrinsic::OverloadFlag;
+using OverloadFlags = tint::core::intrinsic::OverloadFlags;
+using OverloadIndex = tint::core::intrinsic::OverloadIndex;
+using OverloadInfo = tint::core::intrinsic::OverloadInfo;
+using ParameterIndex = tint::core::intrinsic::ParameterIndex;
+using ParameterInfo = tint::core::intrinsic::ParameterInfo;
+using StringStream = tint::StringStream;
+using TemplateNumberIndex = tint::core::intrinsic::TemplateNumberIndex;
+using TemplateNumberInfo = tint::core::intrinsic::TemplateNumberInfo;
+using TemplateTypeIndex = tint::core::intrinsic::TemplateTypeIndex;
+using TemplateTypeInfo = tint::core::intrinsic::TemplateTypeInfo;
+using Type = tint::core::type::Type;
+using TypeMatcher = tint::core::intrinsic::TypeMatcher;
+using TypeMatcherIndex = tint::core::intrinsic::TypeMatcherIndex;
+using TypeMatcherIndicesIndex = tint::core::intrinsic::TypeMatcherIndicesIndex;
+
+template <size_t N>
+using TemplateNumberMatcher = tint::core::intrinsic::TemplateNumberMatcher<N>;
+
+template <size_t N>
+using TemplateTypeMatcher = tint::core::intrinsic::TemplateTypeMatcher<N>;
+
+// clang-format off
+
+/// TypeMatcher for 'type u32'
+constexpr TypeMatcher kU32Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchU32(state, ty)) {
+      return nullptr;
+    }
+    return BuildU32(state, ty);
+  },
+/* string */ [](MatchState*) -> std::string {
+    return "u32";
+  }
+};
+
+
+/// Type and number matchers
+
+/// The template types, types, and type matchers
+constexpr TypeMatcher kTypeMatchers[] = {
+  /* [0] */ kU32Matcher,
+};
+
+constexpr TypeMatcherIndex kTypeMatcherIndices[] = {
+  /* [0] */ TypeMatcherIndex(0),
+};
+
+static_assert(TypeMatcherIndex::CanIndex(kTypeMatcherIndices),
+              "TypeMatcherIndex is not large enough to index kTypeMatcherIndices");
+
+constexpr ParameterInfo kParameters[] = {
+  {
+    /* [0] */
+    /* usage */ core::ParameterUsage::kNone,
+    /* type_matcher_indices */ TypeMatcherIndicesIndex(0),
+    /* number_matcher_indices */ NumberMatcherIndicesIndex(/* invalid */),
+  },
+};
+
+static_assert(ParameterIndex::CanIndex(kParameters),
+              "ParameterIndex is not large enough to index kParameters");
+
+constexpr OverloadInfo kOverloads[] = {
+  {
+    /* [0] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 1,
+    /* num_template_types */ 0,
+    /* num_template_numbers */ 0,
+    /* template_types */ TemplateTypeIndex(/* invalid */),
+    /* template_numbers */ TemplateNumberIndex(/* invalid */),
+    /* parameters */ ParameterIndex(0),
+    /* return_type_matcher_indices */ TypeMatcherIndicesIndex(/* invalid */),
+    /* return_number_matcher_indices */ NumberMatcherIndicesIndex(/* invalid */),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+};
+
+static_assert(OverloadIndex::CanIndex(kOverloads),
+              "OverloadIndex is not large enough to index kOverloads");
+
+constexpr IntrinsicInfo kBuiltins[] = {
+  {
+    /* [0] */
+    /* fn threadgroup_barrier(u32) */
+    /* num overloads */ 1,
+    /* overloads */ OverloadIndex(0),
+  },
+};
+
+// clang-format on
+
+}  // anonymous namespace
+
+const core::intrinsic::TableData Dialect::kData{
+    /* template_types */ Empty,
+    /* template_numbers */ Empty,
+    /* type_matcher_indices */ kTypeMatcherIndices,
+    /* number_matcher_indices */ Empty,
+    /* type_matchers */ kTypeMatchers,
+    /* number_matchers */ Empty,
+    /* parameters */ kParameters,
+    /* overloads */ kOverloads,
+    /* const_eval_functions */ Empty,
+    /* ctor_conv */ Empty,
+    /* builtins */ kBuiltins,
+    /* binary '+' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '-' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '*' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '/' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '%' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '^' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '&' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '|' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '&&' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '||' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '==' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '!=' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '<' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '>' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '<=' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '>=' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '<<' */ tint::core::intrinsic::kNoOverloads,
+    /* binary '>>' */ tint::core::intrinsic::kNoOverloads,
+    /* unary '!' */ tint::core::intrinsic::kNoOverloads,
+    /* unary '~' */ tint::core::intrinsic::kNoOverloads,
+    /* unary '-' */ tint::core::intrinsic::kNoOverloads,
+};
+
+}  // namespace tint::msl::intrinsic
diff --git a/src/tint/lang/msl/intrinsic/data.cc.tmpl b/src/tint/lang/msl/intrinsic/data.cc.tmpl
new file mode 100644
index 0000000..2314b5a
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/data.cc.tmpl
@@ -0,0 +1,34 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate intrinsic_table.inl
+Used by BuiltinTable.cc for builtin overload resolution.
+
+To update the generated file, run:
+    ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- Import "src/tint/utils/templates/intrinsic_table_data.tmpl.inc" -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+
+#include <limits>
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/type_matchers.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/text/string_stream.h"
+
+namespace tint::msl::intrinsic {
+
+using namespace tint::core::intrinsic;  // NOLINT(build/namespaces)
+
+{{ Eval "Data"
+  "Intrinsics" $I
+  "Name"       "Dialect::kData" -}}
+
+}  // namespace tint::msl::intrinsic
diff --git a/src/tint/lang/msl/intrinsic/dialect.h b/src/tint/lang/msl/intrinsic/dialect.h
new file mode 100644
index 0000000..700df35
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/dialect.h
@@ -0,0 +1,51 @@
+// Copyright 2023 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_MSL_INTRINSIC_DIALECT_H_
+#define SRC_TINT_LANG_MSL_INTRINSIC_DIALECT_H_
+
+#include "src/tint/lang/core/intrinsic/table_data.h"
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl::intrinsic {
+
+/// Dialect holds the intrinsic table data and types for the MSL dialect
+struct Dialect {
+    /// The dialect's intrinsic table data
+    static const core::intrinsic::TableData kData;
+
+    /// The dialect's builtin function enumerator
+    using BuiltinFn = msl::BuiltinFn;
+
+    /// @returns the name of the builtin function @p fn
+    /// @param fn the builtin function
+    static std::string_view ToString(BuiltinFn fn) { return str(fn); }
+};
+
+}  // namespace tint::msl::intrinsic
+
+#endif  // SRC_TINT_LANG_MSL_INTRINSIC_DIALECT_H_
diff --git a/src/tint/lang/msl/ir/BUILD.bazel b/src/tint/lang/msl/ir/BUILD.bazel
new file mode 100644
index 0000000..cff7707
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.bazel
@@ -0,0 +1,108 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.bazel.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+load("//src/tint:flags.bzl", "COPTS")
+load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+  name = "ir",
+  srcs = [
+    "builtin_call.cc",
+  ],
+  hdrs = [
+    "builtin_call.h",
+  ],
+  deps = [
+    "//src/tint/api/common",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/ir",
+    "//src/tint/lang/core/type",
+    "//src/tint/lang/msl",
+    "//src/tint/lang/msl/intrinsic",
+    "//src/tint/utils/containers",
+    "//src/tint/utils/diagnostic",
+    "//src/tint/utils/ice",
+    "//src/tint/utils/id",
+    "//src/tint/utils/macros",
+    "//src/tint/utils/math",
+    "//src/tint/utils/memory",
+    "//src/tint/utils/reflection",
+    "//src/tint/utils/result",
+    "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
+    "//src/tint/utils/text",
+    "//src/tint/utils/traits",
+  ],
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
+cc_library(
+  name = "test",
+  alwayslink = True,
+  srcs = [
+    "builtin_call_test.cc",
+  ],
+  deps = [
+    "//src/tint/api/common",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/ir",
+    "//src/tint/lang/core/ir:test",
+    "//src/tint/lang/core/type",
+    "//src/tint/lang/msl",
+    "//src/tint/lang/msl/intrinsic",
+    "//src/tint/lang/msl/ir",
+    "//src/tint/utils/containers",
+    "//src/tint/utils/diagnostic",
+    "//src/tint/utils/ice",
+    "//src/tint/utils/id",
+    "//src/tint/utils/macros",
+    "//src/tint/utils/math",
+    "//src/tint/utils/memory",
+    "//src/tint/utils/reflection",
+    "//src/tint/utils/result",
+    "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
+    "//src/tint/utils/text",
+    "//src/tint/utils/traits",
+    "@gtest",
+  ],
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/msl/ir/BUILD.cmake b/src/tint/lang/msl/ir/BUILD.cmake
new file mode 100644
index 0000000..1f4769d
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.cmake
@@ -0,0 +1,106 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.cmake.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+################################################################################
+# Target:    tint_lang_msl_ir
+# Kind:      lib
+################################################################################
+tint_add_target(tint_lang_msl_ir lib
+  lang/msl/ir/builtin_call.cc
+  lang/msl/ir/builtin_call.h
+)
+
+tint_target_add_dependencies(tint_lang_msl_ir lib
+  tint_api_common
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_ir
+  tint_lang_core_type
+  tint_lang_msl
+  tint_lang_msl_intrinsic
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+################################################################################
+# Target:    tint_lang_msl_ir_test
+# Kind:      test
+################################################################################
+tint_add_target(tint_lang_msl_ir_test test
+  lang/msl/ir/builtin_call_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_msl_ir_test test
+  tint_api_common
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_ir
+  tint_lang_core_ir_test
+  tint_lang_core_type
+  tint_lang_msl
+  tint_lang_msl_intrinsic
+  tint_lang_msl_ir
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_msl_ir_test test
+  "gtest"
+)
diff --git a/src/tint/lang/msl/ir/BUILD.gn b/src/tint/lang/msl/ir/BUILD.gn
new file mode 100644
index 0000000..a072d1c
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.gn
@@ -0,0 +1,104 @@
+# Copyright 2023 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+#   tools/src/cmd/gen/build/BUILD.gn.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+#                       Do not modify this file directly
+################################################################################
+
+import("../../../../../scripts/tint_overrides_with_defaults.gni")
+
+import("${tint_src_dir}/tint.gni")
+
+if (tint_build_unittests || tint_build_benchmarks) {
+  import("//testing/test.gni")
+}
+
+libtint_source_set("ir") {
+  sources = [
+    "builtin_call.cc",
+    "builtin_call.h",
+  ]
+  deps = [
+    "${tint_src_dir}/api/common",
+    "${tint_src_dir}/lang/core",
+    "${tint_src_dir}/lang/core/constant",
+    "${tint_src_dir}/lang/core/intrinsic",
+    "${tint_src_dir}/lang/core/ir",
+    "${tint_src_dir}/lang/core/type",
+    "${tint_src_dir}/lang/msl",
+    "${tint_src_dir}/lang/msl/intrinsic",
+    "${tint_src_dir}/utils/containers",
+    "${tint_src_dir}/utils/diagnostic",
+    "${tint_src_dir}/utils/ice",
+    "${tint_src_dir}/utils/id",
+    "${tint_src_dir}/utils/macros",
+    "${tint_src_dir}/utils/math",
+    "${tint_src_dir}/utils/memory",
+    "${tint_src_dir}/utils/reflection",
+    "${tint_src_dir}/utils/result",
+    "${tint_src_dir}/utils/rtti",
+    "${tint_src_dir}/utils/symbol",
+    "${tint_src_dir}/utils/text",
+    "${tint_src_dir}/utils/traits",
+  ]
+}
+if (tint_build_unittests) {
+  tint_unittests_source_set("unittests") {
+    sources = [ "builtin_call_test.cc" ]
+    deps = [
+      "${tint_src_dir}:gmock_and_gtest",
+      "${tint_src_dir}/api/common",
+      "${tint_src_dir}/lang/core",
+      "${tint_src_dir}/lang/core/constant",
+      "${tint_src_dir}/lang/core/intrinsic",
+      "${tint_src_dir}/lang/core/ir",
+      "${tint_src_dir}/lang/core/ir:unittests",
+      "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/msl",
+      "${tint_src_dir}/lang/msl/intrinsic",
+      "${tint_src_dir}/lang/msl/ir",
+      "${tint_src_dir}/utils/containers",
+      "${tint_src_dir}/utils/diagnostic",
+      "${tint_src_dir}/utils/ice",
+      "${tint_src_dir}/utils/id",
+      "${tint_src_dir}/utils/macros",
+      "${tint_src_dir}/utils/math",
+      "${tint_src_dir}/utils/memory",
+      "${tint_src_dir}/utils/reflection",
+      "${tint_src_dir}/utils/result",
+      "${tint_src_dir}/utils/rtti",
+      "${tint_src_dir}/utils/symbol",
+      "${tint_src_dir}/utils/text",
+      "${tint_src_dir}/utils/traits",
+    ]
+  }
+}
diff --git a/src/tint/lang/msl/ir/builtin_call.cc b/src/tint/lang/msl/ir/builtin_call.cc
new file mode 100644
index 0000000..8e20da3
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call.cc
@@ -0,0 +1,56 @@
+// Copyright 2023 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/msl/ir/builtin_call.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/utils/ice/ice.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::msl::ir::BuiltinCall);
+
+namespace tint::msl::ir {
+
+BuiltinCall::BuiltinCall(core::ir::InstructionResult* result,
+                         BuiltinFn func,
+                         VectorRef<core::ir::Value*> arguments)
+    : Base(result, arguments), func_(func) {
+    flags_.Add(Flag::kSequenced);
+    TINT_ASSERT(func != BuiltinFn::kNone);
+}
+
+BuiltinCall::~BuiltinCall() = default;
+
+BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
+    auto* new_result = ctx.Clone(Result(0));
+    auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
+    return ctx.ir.instructions.Create<BuiltinCall>(new_result, func_, new_args);
+}
+
+}  // namespace tint::msl::ir
diff --git a/src/tint/lang/msl/ir/builtin_call.h b/src/tint/lang/msl/ir/builtin_call.h
new file mode 100644
index 0000000..de63d4d
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call.h
@@ -0,0 +1,76 @@
+// Copyright 2023 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_MSL_IR_BUILTIN_CALL_H_
+#define SRC_TINT_LANG_MSL_IR_BUILTIN_CALL_H_
+
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/table_data.h"
+#include "src/tint/lang/core/ir/builtin_call.h"
+#include "src/tint/lang/msl/builtin_fn.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/rtti/castable.h"
+
+namespace tint::msl::ir {
+
+/// A msl builtin call instruction in the IR.
+class BuiltinCall final : public Castable<BuiltinCall, core::ir::BuiltinCall> {
+  public:
+    /// Constructor
+    /// @param result the result value
+    /// @param func the builtin function
+    /// @param args the conversion arguments
+    BuiltinCall(core::ir::InstructionResult* result,
+                BuiltinFn func,
+                VectorRef<core::ir::Value*> args = tint::Empty);
+    ~BuiltinCall() override;
+
+    /// @copydoc core::ir::Instruction::Clone()
+    BuiltinCall* Clone(core::ir::CloneContext& ctx) override;
+
+    /// @returns the builtin function
+    BuiltinFn Func() const { return func_; }
+
+    /// @returns the identifier for the function
+    size_t FuncId() const override { return static_cast<size_t>(func_); }
+
+    /// @returns the friendly name for the instruction
+    std::string FriendlyName() const override { return std::string("msl.") + str(func_); }
+
+    /// @returns the table data to validate this builtin
+    const core::intrinsic::TableData& TableData() const override {
+        return msl::intrinsic::Dialect::kData;
+    }
+
+  private:
+    BuiltinFn func_;
+};
+
+}  // namespace tint::msl::ir
+
+#endif  // SRC_TINT_LANG_MSL_IR_BUILTIN_CALL_H_
diff --git a/src/tint/lang/msl/ir/builtin_call_test.cc b/src/tint/lang/msl/ir/builtin_call_test.cc
new file mode 100644
index 0000000..cbe47ba
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call_test.cc
@@ -0,0 +1,59 @@
+// Copyright 2023 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/msl/ir/builtin_call.h"
+#include "gmock/gmock.h"
+#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/ir_helper_test.h"
+
+namespace tint::msl::ir {
+namespace {
+
+using namespace tint::core::number_suffixes;  // NOLINT
+                                              //
+using IR_MslBuiltinCallTest = core::ir::IRTestHelper;
+
+TEST_F(IR_MslBuiltinCallTest, Clone) {
+    auto* builtin = b.Call<BuiltinCall>(mod.Types().void_(), BuiltinFn::kThreadgroupBarrier, 0_u);
+
+    auto* new_b = clone_ctx.Clone(builtin);
+
+    EXPECT_NE(builtin, new_b);
+    EXPECT_NE(builtin->Result(0), new_b->Result(0));
+    EXPECT_EQ(mod.Types().void_(), new_b->Result(0)->Type());
+
+    EXPECT_EQ(BuiltinFn::kThreadgroupBarrier, new_b->Func());
+
+    auto args = new_b->Args();
+    EXPECT_EQ(1u, args.Length());
+
+    auto* val0 = args[0]->As<core::ir::Constant>()->Value();
+    EXPECT_EQ(0_u, val0->As<core::constant::Scalar<core::u32>>()->ValueAs<core::u32>());
+}
+
+}  // namespace
+}  // namespace tint::msl::ir
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
new file mode 100644
index 0000000..7709699
--- /dev/null
+++ b/src/tint/lang/msl/msl.def
@@ -0,0 +1,38 @@
+// Copyright 2023 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// MSL builtin definition file                                                //
+////////////////////////////////////////////////////////////////////////////////
+
+type u32
+
+////////////////////////////////////////////////////////////////////////////////
+// Builtin Functions                                                          //
+////////////////////////////////////////////////////////////////////////////////
+@stage("compute") fn threadgroup_barrier(u32)
+
diff --git a/src/tint/lang/msl/writer/printer/BUILD.bazel b/src/tint/lang/msl/writer/printer/BUILD.bazel
index 7de0377..ef25cfe 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/printer/BUILD.bazel
@@ -51,6 +51,9 @@
     "//src/tint/lang/core/intrinsic",
     "//src/tint/lang/core/ir",
     "//src/tint/lang/core/type",
+    "//src/tint/lang/msl",
+    "//src/tint/lang/msl/intrinsic",
+    "//src/tint/lang/msl/ir",
     "//src/tint/utils/containers",
     "//src/tint/utils/diagnostic",
     "//src/tint/utils/generator",
diff --git a/src/tint/lang/msl/writer/printer/BUILD.cmake b/src/tint/lang/msl/writer/printer/BUILD.cmake
index 37d0fcb..b922ea6 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/printer/BUILD.cmake
@@ -52,6 +52,9 @@
   tint_lang_core_intrinsic
   tint_lang_core_ir
   tint_lang_core_type
+  tint_lang_msl
+  tint_lang_msl_intrinsic
+  tint_lang_msl_ir
   tint_utils_containers
   tint_utils_diagnostic
   tint_utils_generator
diff --git a/src/tint/lang/msl/writer/printer/BUILD.gn b/src/tint/lang/msl/writer/printer/BUILD.gn
index ce460cc..3db1a8f 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.gn
+++ b/src/tint/lang/msl/writer/printer/BUILD.gn
@@ -54,6 +54,9 @@
       "${tint_src_dir}/lang/core/intrinsic",
       "${tint_src_dir}/lang/core/ir",
       "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/msl",
+      "${tint_src_dir}/lang/msl/intrinsic",
+      "${tint_src_dir}/lang/msl/ir",
       "${tint_src_dir}/utils/containers",
       "${tint_src_dir}/utils/diagnostic",
       "${tint_src_dir}/utils/generator",
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 2ededba..24b4b29 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -83,6 +83,8 @@
 #include "src/tint/lang/core/type/u32.h"
 #include "src/tint/lang/core/type/vector.h"
 #include "src/tint/lang/core/type/void.h"
+#include "src/tint/lang/msl/barrier_type.h"
+#include "src/tint/lang/msl/ir/builtin_call.h"
 #include "src/tint/lang/msl/writer/common/printer_support.h"
 #include "src/tint/utils/containers/map.h"
 #include "src/tint/utils/generator/text_generator.h"
@@ -372,6 +374,7 @@
                     [&](const core::ir::Var* var) { out << NameOf(var->Result(0)); },          //
                     [&](const core::ir::Bitcast* b) { EmitBitcast(out, b); },                  //
                     [&](const core::ir::Access* a) { EmitAccess(out, a); },                    //
+                    [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); },        //
                     [&](const core::ir::CoreBuiltinCall* c) { EmitCoreBuiltinCall(out, c); },  //
                     [&](const core::ir::UserCall* c) { EmitUserCall(out, c); },                //
                     [&](const core::ir::LoadVectorElement* e) {
@@ -753,6 +756,37 @@
         }
     }
 
+    void EmitMslBuiltinCall(StringStream& out, const msl::ir::BuiltinCall* c) {
+        switch (c->Func()) {
+            case msl::BuiltinFn::kThreadgroupBarrier: {
+                auto flags = c->Args()[0]->As<core::ir::Constant>()->Value()->ValueAs<uint8_t>();
+                out << "threadgroup_barrier(";
+                bool emitted_flag = false;
+
+                auto emit = [&](BarrierType type, const std::string& name) {
+                    if ((flags & type) != type) {
+                        return;
+                    }
+
+                    if (emitted_flag) {
+                        out << " | ";
+                    }
+                    emitted_flag = true;
+                    out << "mem_flags::mem_" << name;
+                };
+                emit(BarrierType::kDevice, "device");
+                emit(BarrierType::kThreadGroup, "threadgroup");
+                emit(BarrierType::kTexture, "texture");
+
+                out << ")";
+                return;
+            }
+            default:
+                TINT_ICE() << "undefined MSL ir function";
+                return;
+        }
+    }
+
     void EmitCoreBuiltinCall(StringStream& out, const core::ir::CoreBuiltinCall* c) {
         EmitCoreBuiltinName(out, c->Func());
         out << "(";
diff --git a/src/tint/lang/msl/writer/raise/BUILD.bazel b/src/tint/lang/msl/writer/raise/BUILD.bazel
index a192956..c906eaa 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.bazel
+++ b/src/tint/lang/msl/writer/raise/BUILD.bazel
@@ -39,24 +39,36 @@
 cc_library(
   name = "raise",
   srcs = [
+    "builtin_polyfill.cc",
     "raise.cc",
   ],
   hdrs = [
+    "builtin_polyfill.h",
     "raise.h",
   ],
   deps = [
     "//src/tint/api/common",
     "//src/tint/api/options",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/ir",
     "//src/tint/lang/core/ir/transform",
+    "//src/tint/lang/core/type",
+    "//src/tint/lang/msl",
+    "//src/tint/lang/msl/intrinsic",
+    "//src/tint/lang/msl/ir",
     "//src/tint/utils/containers",
     "//src/tint/utils/diagnostic",
     "//src/tint/utils/ice",
+    "//src/tint/utils/id",
     "//src/tint/utils/macros",
     "//src/tint/utils/math",
     "//src/tint/utils/memory",
     "//src/tint/utils/reflection",
     "//src/tint/utils/result",
     "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
     "//src/tint/utils/text",
     "//src/tint/utils/traits",
   ] + select({
@@ -68,6 +80,43 @@
   copts = COPTS,
   visibility = ["//visibility:public"],
 )
+cc_library(
+  name = "test",
+  alwayslink = True,
+  srcs = [
+    "builtin_polyfill_test.cc",
+  ],
+  deps = [
+    "//src/tint/api/common",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/ir",
+    "//src/tint/lang/core/ir/transform:test",
+    "//src/tint/lang/core/type",
+    "//src/tint/utils/containers",
+    "//src/tint/utils/diagnostic",
+    "//src/tint/utils/ice",
+    "//src/tint/utils/id",
+    "//src/tint/utils/macros",
+    "//src/tint/utils/math",
+    "//src/tint/utils/memory",
+    "//src/tint/utils/reflection",
+    "//src/tint/utils/result",
+    "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
+    "//src/tint/utils/text",
+    "//src/tint/utils/traits",
+    "@gtest",
+  ] + select({
+    ":tint_build_msl_writer": [
+      "//src/tint/lang/msl/writer/raise",
+    ],
+    "//conditions:default": [],
+  }),
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
 
 alias(
   name = "tint_build_msl_writer",
diff --git a/src/tint/lang/msl/writer/raise/BUILD.cmake b/src/tint/lang/msl/writer/raise/BUILD.cmake
index 0e41a3c..06e15b9 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/msl/writer/raise/BUILD.cmake
@@ -41,6 +41,8 @@
 # Condition: TINT_BUILD_MSL_WRITER
 ################################################################################
 tint_add_target(tint_lang_msl_writer_raise lib
+  lang/msl/writer/raise/builtin_polyfill.cc
+  lang/msl/writer/raise/builtin_polyfill.h
   lang/msl/writer/raise/raise.cc
   lang/msl/writer/raise/raise.h
 )
@@ -48,16 +50,26 @@
 tint_target_add_dependencies(tint_lang_msl_writer_raise lib
   tint_api_common
   tint_api_options
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_ir
   tint_lang_core_ir_transform
+  tint_lang_core_type
+  tint_lang_msl
+  tint_lang_msl_intrinsic
+  tint_lang_msl_ir
   tint_utils_containers
   tint_utils_diagnostic
   tint_utils_ice
+  tint_utils_id
   tint_utils_macros
   tint_utils_math
   tint_utils_memory
   tint_utils_reflection
   tint_utils_result
   tint_utils_rtti
+  tint_utils_symbol
   tint_utils_text
   tint_utils_traits
 )
@@ -68,4 +80,48 @@
   )
 endif(TINT_BUILD_MSL_WRITER)
 
+endif(TINT_BUILD_MSL_WRITER)
+if(TINT_BUILD_MSL_WRITER)
+################################################################################
+# Target:    tint_lang_msl_writer_raise_test
+# Kind:      test
+# Condition: TINT_BUILD_MSL_WRITER
+################################################################################
+tint_add_target(tint_lang_msl_writer_raise_test test
+  lang/msl/writer/raise/builtin_polyfill_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_msl_writer_raise_test test
+  tint_api_common
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_ir
+  tint_lang_core_ir_transform_test
+  tint_lang_core_type
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_msl_writer_raise_test test
+  "gtest"
+)
+
+if(TINT_BUILD_MSL_WRITER)
+  tint_target_add_dependencies(tint_lang_msl_writer_raise_test test
+    tint_lang_msl_writer_raise
+  )
+endif(TINT_BUILD_MSL_WRITER)
+
 endif(TINT_BUILD_MSL_WRITER)
\ No newline at end of file
diff --git a/src/tint/lang/msl/writer/raise/BUILD.gn b/src/tint/lang/msl/writer/raise/BUILD.gn
index 74cc477..d8e6ad5 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.gn
+++ b/src/tint/lang/msl/writer/raise/BUILD.gn
@@ -37,25 +37,41 @@
 import("../../../../../../scripts/tint_overrides_with_defaults.gni")
 
 import("${tint_src_dir}/tint.gni")
+
+if (tint_build_unittests || tint_build_benchmarks) {
+  import("//testing/test.gni")
+}
 if (tint_build_msl_writer) {
   libtint_source_set("raise") {
     sources = [
+      "builtin_polyfill.cc",
+      "builtin_polyfill.h",
       "raise.cc",
       "raise.h",
     ]
     deps = [
       "${tint_src_dir}/api/common",
       "${tint_src_dir}/api/options",
+      "${tint_src_dir}/lang/core",
+      "${tint_src_dir}/lang/core/constant",
+      "${tint_src_dir}/lang/core/intrinsic",
+      "${tint_src_dir}/lang/core/ir",
       "${tint_src_dir}/lang/core/ir/transform",
+      "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/msl",
+      "${tint_src_dir}/lang/msl/intrinsic",
+      "${tint_src_dir}/lang/msl/ir",
       "${tint_src_dir}/utils/containers",
       "${tint_src_dir}/utils/diagnostic",
       "${tint_src_dir}/utils/ice",
+      "${tint_src_dir}/utils/id",
       "${tint_src_dir}/utils/macros",
       "${tint_src_dir}/utils/math",
       "${tint_src_dir}/utils/memory",
       "${tint_src_dir}/utils/reflection",
       "${tint_src_dir}/utils/result",
       "${tint_src_dir}/utils/rtti",
+      "${tint_src_dir}/utils/symbol",
       "${tint_src_dir}/utils/text",
       "${tint_src_dir}/utils/traits",
     ]
@@ -65,3 +81,37 @@
     }
   }
 }
+if (tint_build_unittests) {
+  if (tint_build_msl_writer) {
+    tint_unittests_source_set("unittests") {
+      sources = [ "builtin_polyfill_test.cc" ]
+      deps = [
+        "${tint_src_dir}:gmock_and_gtest",
+        "${tint_src_dir}/api/common",
+        "${tint_src_dir}/lang/core",
+        "${tint_src_dir}/lang/core/constant",
+        "${tint_src_dir}/lang/core/intrinsic",
+        "${tint_src_dir}/lang/core/ir",
+        "${tint_src_dir}/lang/core/ir/transform:unittests",
+        "${tint_src_dir}/lang/core/type",
+        "${tint_src_dir}/utils/containers",
+        "${tint_src_dir}/utils/diagnostic",
+        "${tint_src_dir}/utils/ice",
+        "${tint_src_dir}/utils/id",
+        "${tint_src_dir}/utils/macros",
+        "${tint_src_dir}/utils/math",
+        "${tint_src_dir}/utils/memory",
+        "${tint_src_dir}/utils/reflection",
+        "${tint_src_dir}/utils/result",
+        "${tint_src_dir}/utils/rtti",
+        "${tint_src_dir}/utils/symbol",
+        "${tint_src_dir}/utils/text",
+        "${tint_src_dir}/utils/traits",
+      ]
+
+      if (tint_build_msl_writer) {
+        deps += [ "${tint_src_dir}/lang/msl/writer/raise" ]
+      }
+    }
+  }
+}
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
new file mode 100644
index 0000000..1905246
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -0,0 +1,157 @@
+// Copyright 2023 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/msl/writer/raise/builtin_polyfill.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/constant.h"
+#include "src/tint/lang/core/ir/core_builtin_call.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/msl/barrier_type.h"
+#include "src/tint/lang/msl/ir/builtin_call.h"
+
+namespace tint::msl::writer::raise {
+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() {
+        // Find the builtins that need replacing.
+        Vector<core::ir::CoreBuiltinCall*, 4> worklist;
+        for (auto* inst : ir.instructions.Objects()) {
+            if (!inst->Alive()) {
+                continue;
+            }
+            if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
+                switch (builtin->Func()) {
+                    case core::BuiltinFn::kStorageBarrier:
+                    case core::BuiltinFn::kWorkgroupBarrier:
+                    case core::BuiltinFn::kTextureBarrier:
+                        worklist.Push(builtin);
+                        break;
+                    default:
+                        break;
+                }
+            }
+        }
+
+        // Replace the builtins that we found.
+        for (auto* builtin : worklist) {
+            core::ir::Value* replacement = nullptr;
+            switch (builtin->Func()) {
+                case core::BuiltinFn::kStorageBarrier:
+                    replacement = StorageBarrier(builtin);
+                    break;
+                case core::BuiltinFn::kWorkgroupBarrier:
+                    replacement = WorkgroupBarrier(builtin);
+                    break;
+                case core::BuiltinFn::kTextureBarrier:
+                    replacement = TextureBarrier(builtin);
+                    break;
+                default:
+                    break;
+            }
+            TINT_ASSERT_OR_RETURN(replacement);
+
+            // Replace the old builtin result with the new value.
+            if (auto name = ir.NameOf(builtin->Result(0))) {
+                ir.SetName(replacement, name);
+            }
+            builtin->Result(0)->ReplaceAllUsesWith(replacement);
+            builtin->Destroy();
+        }
+    }
+
+    /// Handle a `workgroupBarrier()` builtin.
+    /// @param builtin the builtin call instruction
+    /// @returns the replacement value
+    core::ir::Value* WorkgroupBarrier(core::ir::CoreBuiltinCall* builtin) {
+        // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+        auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kThreadGroup))};
+
+        auto* call = b.Call<msl::ir::BuiltinCall>(
+            builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+        call->InsertBefore(builtin);
+        return call->Result(0);
+    }
+
+    /// Handle a `storageBarrier()` builtin.
+    /// @param builtin the builtin call instruction
+    /// @returns the replacement value
+    core::ir::Value* StorageBarrier(core::ir::CoreBuiltinCall* builtin) {
+        // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+        auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kDevice))};
+
+        auto* call = b.Call<msl::ir::BuiltinCall>(
+            builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+        call->InsertBefore(builtin);
+        return call->Result(0);
+    }
+
+    /// Handle a `textureBarrier()` builtin.
+    /// @param builtin the builtin call instruction
+    /// @returns the replacement value
+    core::ir::Value* TextureBarrier(core::ir::CoreBuiltinCall* builtin) {
+        // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+        auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kTexture))};
+
+        auto* call = b.Call<msl::ir::BuiltinCall>(
+            builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+        call->InsertBefore(builtin);
+        return call->Result(0);
+    }
+};
+
+}  // namespace
+
+Result<SuccessType> BuiltinPolyfill(core::ir::Module& ir) {
+    auto result = ValidateAndDumpIfNeeded(ir, "BuiltinPolyfill transform");
+    if (!result) {
+        return result.Failure();
+    }
+
+    State{ir}.Process();
+
+    return Success;
+}
+
+}  // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.h b/src/tint/lang/msl/writer/raise/builtin_polyfill.h
new file mode 100644
index 0000000..344ed14
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.h
@@ -0,0 +1,54 @@
+// Copyright 2023 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_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
+#define SRC_TINT_LANG_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
+
+#include <string>
+
+#include "src/tint/lang/core/ir/constant.h"
+#include "src/tint/lang/core/type/type.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+class Texture;
+}  // namespace tint::core::ir
+
+namespace tint::msl::writer::raise {
+
+/// BuiltinPolyfill is a transform that replaces calls to builtins with polyfills and calls to
+/// MSL backend intrinsic functions.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> BuiltinPolyfill(core::ir::Module& module);
+
+}  // namespace tint::msl::writer::raise
+
+#endif  // SRC_TINT_LANG_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
new file mode 100644
index 0000000..b61a9a0
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -0,0 +1,143 @@
+// Copyright 2023 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/msl/writer/raise/builtin_polyfill.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/type/array.h"
+#include "src/tint/lang/core/type/atomic.h"
+#include "src/tint/lang/core/type/builtin_structs.h"
+#include "src/tint/lang/core/type/depth_texture.h"
+#include "src/tint/lang/core/type/multisampled_texture.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
+#include "src/tint/lang/core/type/storage_texture.h"
+
+namespace tint::msl::writer::raise {
+namespace {
+
+using MslWriter_BuiltinPolyfillTest = core::ir::transform::TransformTest;
+
+TEST_F(MslWriter_BuiltinPolyfillTest, WorkgroupBarrier) {
+    auto* func = b.Function("foo", ty.void_());
+    func->SetStage(core::ir::Function::PipelineStage::kCompute);
+    b.Append(func->Block(), [&] {
+        b.Call(ty.void_(), core::BuiltinFn::kWorkgroupBarrier);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = workgroupBarrier
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = msl.threadgroup_barrier 4u
+    ret
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, StorageBarrier) {
+    auto* func = b.Function("foo", ty.void_());
+    func->SetStage(core::ir::Function::PipelineStage::kCompute);
+    b.Append(func->Block(), [&] {
+        b.Call(ty.void_(), core::BuiltinFn::kStorageBarrier);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = storageBarrier
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = msl.threadgroup_barrier 1u
+    ret
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, TextureBarrier) {
+    auto* func = b.Function("foo", ty.void_());
+    func->SetStage(core::ir::Function::PipelineStage::kCompute);
+    b.Append(func->Block(), [&] {
+        b.Call(ty.void_(), core::BuiltinFn::kTextureBarrier);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = textureBarrier
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+  %b1 = block {
+    %2:void = msl.threadgroup_barrier 2u
+    ret
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
+}  // namespace
+}  // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index da70216..4140f46 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -41,6 +41,7 @@
 #include "src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.h"
 #include "src/tint/lang/core/ir/transform/zero_init_workgroup_memory.h"
 #include "src/tint/lang/msl/writer/common/option_helpers.h"
+#include "src/tint/lang/msl/writer/raise/builtin_polyfill.h"
 
 namespace tint::msl::writer::raise {
 
@@ -96,13 +97,13 @@
 
     // PreservePadding must come before DirectVariableAccess.
     RUN_TRANSFORM(core::ir::transform::PreservePadding);
-
     RUN_TRANSFORM(core::ir::transform::VectorizeScalarMatrixConstructors);
 
     // DemoteToHelper must come before any transform that introduces non-core instructions.
     RUN_TRANSFORM(core::ir::transform::DemoteToHelper);
 
     RUN_TRANSFORM(core::ir::transform::ValueToLet);
+    RUN_TRANSFORM(BuiltinPolyfill);
 
     return Success;
 }