diff --git a/src/tint/cmd/loopy/BUILD.bazel b/src/tint/cmd/loopy/BUILD.bazel
index 4498c58..6054294 100644
--- a/src/tint/cmd/loopy/BUILD.bazel
+++ b/src/tint/cmd/loopy/BUILD.bazel
@@ -87,6 +87,7 @@
     ":tint_build_msl_writer": [
       "//src/tint/lang/msl/writer",
       "//src/tint/lang/msl/writer/common",
+      "//src/tint/lang/msl/writer/helpers",
     ],
     "//conditions:default": [],
   }) + select({
diff --git a/src/tint/cmd/loopy/BUILD.cmake b/src/tint/cmd/loopy/BUILD.cmake
index 332536d..af48744 100644
--- a/src/tint/cmd/loopy/BUILD.cmake
+++ b/src/tint/cmd/loopy/BUILD.cmake
@@ -92,6 +92,7 @@
   tint_target_add_dependencies(tint_cmd_loopy_cmd cmd
     tint_lang_msl_writer
     tint_lang_msl_writer_common
+    tint_lang_msl_writer_helpers
   )
 endif(TINT_BUILD_MSL_WRITER)
 
diff --git a/src/tint/cmd/loopy/BUILD.gn b/src/tint/cmd/loopy/BUILD.gn
index e68fff0..3e7fa3d 100644
--- a/src/tint/cmd/loopy/BUILD.gn
+++ b/src/tint/cmd/loopy/BUILD.gn
@@ -89,6 +89,7 @@
     deps += [
       "${tint_src_dir}/lang/msl/writer",
       "${tint_src_dir}/lang/msl/writer/common",
+      "${tint_src_dir}/lang/msl/writer/helpers",
     ]
   }
 
diff --git a/src/tint/cmd/loopy/main.cc b/src/tint/cmd/loopy/main.cc
index e3d789e..6cc2317 100644
--- a/src/tint/cmd/loopy/main.cc
+++ b/src/tint/cmd/loopy/main.cc
@@ -41,6 +41,7 @@
 #endif  // TINT_BUILD_HLSL_WRITER
 
 #if TINT_BUILD_MSL_WRITER
+#include "src/tint/lang/msl/writer/helpers/generate_bindings.h"
 #include "src/tint/lang/msl/writer/writer.h"
 #endif  // TINT_BUILD_MSL_WRITER
 
@@ -258,8 +259,7 @@
     }
 
     tint::msl::writer::Options gen_options;
-    gen_options.external_texture_options.bindings_map =
-        tint::cmd::GenerateExternalTextureBindings(*input_program);
+    gen_options.bindings = tint::msl::writer::GenerateBindings(program);
     gen_options.array_length_from_uniform.ubo_binding = tint::BindingPoint{0, 30};
     gen_options.array_length_from_uniform.bindpoint_to_size_index.emplace(tint::BindingPoint{0, 0},
                                                                           0);
diff --git a/src/tint/cmd/tint/BUILD.bazel b/src/tint/cmd/tint/BUILD.bazel
index 0b576e2..1abb879 100644
--- a/src/tint/cmd/tint/BUILD.bazel
+++ b/src/tint/cmd/tint/BUILD.bazel
@@ -98,6 +98,7 @@
       "//src/tint/lang/msl/validate",
       "//src/tint/lang/msl/writer",
       "//src/tint/lang/msl/writer/common",
+      "//src/tint/lang/msl/writer/helpers",
     ],
     "//conditions:default": [],
   }) + select({
diff --git a/src/tint/cmd/tint/BUILD.cmake b/src/tint/cmd/tint/BUILD.cmake
index fa5a02e..bd4e4ed 100644
--- a/src/tint/cmd/tint/BUILD.cmake
+++ b/src/tint/cmd/tint/BUILD.cmake
@@ -104,6 +104,7 @@
     tint_lang_msl_validate
     tint_lang_msl_writer
     tint_lang_msl_writer_common
+    tint_lang_msl_writer_helpers
   )
 endif(TINT_BUILD_MSL_WRITER)
 
diff --git a/src/tint/cmd/tint/BUILD.gn b/src/tint/cmd/tint/BUILD.gn
index 02c9071..4bb12ce 100644
--- a/src/tint/cmd/tint/BUILD.gn
+++ b/src/tint/cmd/tint/BUILD.gn
@@ -101,6 +101,7 @@
       "${tint_src_dir}/lang/msl/validate",
       "${tint_src_dir}/lang/msl/writer",
       "${tint_src_dir}/lang/msl/writer/common",
+      "${tint_src_dir}/lang/msl/writer/helpers",
     ]
   }
 
diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index 19a9c85..7ae374d 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc
@@ -83,6 +83,7 @@
 
 #if TINT_BUILD_MSL_WRITER
 #include "src/tint/lang/msl/validate/validate.h"
+#include "src/tint/lang/msl/writer/helpers/generate_bindings.h"
 #include "src/tint/lang/msl/writer/writer.h"
 #endif  // TINT_BUILD_MSL_WRITER
 
@@ -704,8 +705,7 @@
     gen_options.disable_robustness = !options.enable_robustness;
     gen_options.disable_workgroup_init = options.disable_workgroup_init;
     gen_options.pixel_local_options = options.pixel_local_options;
-    gen_options.external_texture_options.bindings_map =
-        tint::cmd::GenerateExternalTextureBindings(*input_program);
+    gen_options.bindings = tint::msl::writer::GenerateBindings(*input_program);
     gen_options.array_length_from_uniform.ubo_binding = tint::BindingPoint{0, 30};
     gen_options.array_length_from_uniform.bindpoint_to_size_index.emplace(tint::BindingPoint{0, 0},
                                                                           0);
diff --git a/src/tint/fuzzers/BUILD.gn b/src/tint/fuzzers/BUILD.gn
index 375d00d..b570ed4 100644
--- a/src/tint/fuzzers/BUILD.gn
+++ b/src/tint/fuzzers/BUILD.gn
@@ -86,6 +86,7 @@
       "${tint_src_dir}/lang/glsl/writer",
       "${tint_src_dir}/lang/hlsl/writer",
       "${tint_src_dir}/lang/msl/writer",
+      "${tint_src_dir}/lang/msl/writer/helpers",
       "${tint_src_dir}/lang/spirv/writer",
       "${tint_src_dir}/lang/spirv/writer/helpers",
       "${tint_src_dir}/lang/wgsl/ast",
diff --git a/src/tint/fuzzers/CMakeLists.txt b/src/tint/fuzzers/CMakeLists.txt
index 3801e7f..136afd3 100644
--- a/src/tint/fuzzers/CMakeLists.txt
+++ b/src/tint/fuzzers/CMakeLists.txt
@@ -51,6 +51,7 @@
     tint_spvheaders_compile_options(${NAME})
     tint_spvtools_compile_options(${NAME})
   endif()
+  target_link_libraries(${NAME} PRIVATE tint_lang_msl_writer_helpers)
   target_link_libraries(${NAME} PRIVATE tint_lang_spirv_writer_helpers)
   target_compile_options(${NAME} PRIVATE -Wno-missing-prototypes)
 endfunction()
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/CMakeLists.txt b/src/tint/fuzzers/tint_ast_fuzzer/CMakeLists.txt
index cdb1d72..fbdb3b9 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/CMakeLists.txt
+++ b/src/tint/fuzzers/tint_ast_fuzzer/CMakeLists.txt
@@ -28,6 +28,7 @@
 function(add_tint_ast_fuzzer NAME)
   add_executable(${NAME} ${NAME}.cc ${AST_FUZZER_SOURCES})
   target_link_libraries(${NAME} PRIVATE libtint_ast_fuzzer tint_lang_spirv_writer_helpers)
+  target_link_libraries(${NAME} PRIVATE libtint_ast_fuzzer tint_lang_msl_writer_helpers)
   tint_fuzzer_compile_options(${NAME})
   if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
     tint_spvheaders_compile_options(${NAME})
diff --git a/src/tint/fuzzers/tint_common_fuzzer.cc b/src/tint/fuzzers/tint_common_fuzzer.cc
index e6ae17b..de7e8f6 100644
--- a/src/tint/fuzzers/tint_common_fuzzer.cc
+++ b/src/tint/fuzzers/tint_common_fuzzer.cc
@@ -58,6 +58,10 @@
 #include "src/tint/lang/spirv/writer/helpers/generate_bindings.h"
 #endif  // TINT_BUILD_SPV_WRITER
 
+#if TINT_BUILD_MSL_WRITER
+#include "src/tint/lang/msl/writer/helpers/generate_bindings.h"
+#endif  // TINT_BUILD_MSL_WRITER
+
 namespace tint::fuzzers {
 
 namespace {
@@ -263,6 +267,9 @@
 
     switch (output_) {
         case OutputFormat::kMSL:
+#if TINT_BUILD_MSL_WRITER
+            options_msl_.bindings = tint::msl::writer::GenerateBindings(program);
+#endif  // TINT_BUILD_MSL_WRITER
             break;
         case OutputFormat::kHLSL:
             break;
@@ -277,7 +284,7 @@
 
     // For the generates which use MultiPlanar, make sure the configuration options are provided so
     // that the transformer will execute.
-    if (output_ == OutputFormat::kMSL || output_ == OutputFormat::kHLSL) {
+    if (output_ == OutputFormat::kHLSL) {
         // Gather external texture binding information
         // Collect next valid binding number per group
         std::unordered_map<uint32_t, uint32_t> group_to_next_binding_number;
@@ -306,7 +313,6 @@
 
         switch (output_) {
             case OutputFormat::kMSL: {
-                options_msl_.external_texture_options.bindings_map = new_bindings_map;
                 break;
             }
             case OutputFormat::kHLSL: {
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/CMakeLists.txt b/src/tint/fuzzers/tint_regex_fuzzer/CMakeLists.txt
index 6d77043..b6f0043 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/CMakeLists.txt
+++ b/src/tint/fuzzers/tint_regex_fuzzer/CMakeLists.txt
@@ -28,6 +28,7 @@
 function(add_tint_regex_fuzzer NAME)
   add_executable(${NAME} ${NAME}.cc ${REGEX_FUZZER_SOURCES})
   target_link_libraries(${NAME} PRIVATE libtint_regex_fuzzer tint_lang_spirv_writer_helpers)
+  target_link_libraries(${NAME} PRIVATE libtint_regex_fuzzer tint_lang_msl_writer_helpers)
   tint_fuzzer_compile_options(${NAME})
   tint_spvtools_compile_options(${NAME})
   target_compile_definitions(${NAME} PRIVATE CUSTOM_MUTATOR)
diff --git a/src/tint/lang/msl/writer/BUILD.bazel b/src/tint/lang/msl/writer/BUILD.bazel
index e887293..f57979e 100644
--- a/src/tint/lang/msl/writer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/BUILD.bazel
@@ -124,6 +124,7 @@
     ":tint_build_msl_writer": [
       "//src/tint/lang/msl/writer",
       "//src/tint/lang/msl/writer/common",
+      "//src/tint/lang/msl/writer/helpers",
     ],
     "//conditions:default": [],
   }),
diff --git a/src/tint/lang/msl/writer/BUILD.cmake b/src/tint/lang/msl/writer/BUILD.cmake
index 974531e..243f93b 100644
--- a/src/tint/lang/msl/writer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/BUILD.cmake
@@ -37,6 +37,7 @@
 include(lang/msl/writer/ast_printer/BUILD.cmake)
 include(lang/msl/writer/ast_raise/BUILD.cmake)
 include(lang/msl/writer/common/BUILD.cmake)
+include(lang/msl/writer/helpers/BUILD.cmake)
 include(lang/msl/writer/printer/BUILD.cmake)
 include(lang/msl/writer/raise/BUILD.cmake)
 
@@ -141,6 +142,7 @@
   tint_target_add_dependencies(tint_lang_msl_writer_bench bench
     tint_lang_msl_writer
     tint_lang_msl_writer_common
+    tint_lang_msl_writer_helpers
   )
 endif(TINT_BUILD_MSL_WRITER)
 
diff --git a/src/tint/lang/msl/writer/BUILD.gn b/src/tint/lang/msl/writer/BUILD.gn
index a52c10e..2b379e5 100644
--- a/src/tint/lang/msl/writer/BUILD.gn
+++ b/src/tint/lang/msl/writer/BUILD.gn
@@ -126,6 +126,7 @@
         deps += [
           "${tint_src_dir}/lang/msl/writer",
           "${tint_src_dir}/lang/msl/writer/common",
+          "${tint_src_dir}/lang/msl/writer/helpers",
         ]
       }
     }
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 3b09a66..60f7f81 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -59,6 +59,7 @@
 #include "src/tint/lang/msl/writer/ast_raise/packed_vec3.h"
 #include "src/tint/lang/msl/writer/ast_raise/pixel_local.h"
 #include "src/tint/lang/msl/writer/ast_raise/subgroup_ballot.h"
+#include "src/tint/lang/msl/writer/common/option_helpers.h"
 #include "src/tint/lang/msl/writer/common/printer_support.h"
 #include "src/tint/lang/wgsl/ast/alias.h"
 #include "src/tint/lang/wgsl/ast/bool_literal_expression.h"
@@ -191,16 +192,20 @@
         manager.Add<ast::transform::BuiltinPolyfill>();
     }
 
-    // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
-    data.Add<ast::transform::MultiplanarExternalTexture::NewBindingPoints>(
-        options.external_texture_options.bindings_map);
-    manager.Add<ast::transform::MultiplanarExternalTexture>();
+    ExternalTextureOptions external_texture_options{};
+    RemapperData remapper_data{};
+    PopulateRemapperAndMultiplanarOptions(options, remapper_data, external_texture_options);
 
-    // BindingRemapper must come after MultiplanarExternalTexture
     manager.Add<ast::transform::BindingRemapper>();
     data.Add<ast::transform::BindingRemapper::Remappings>(
-        options.binding_remapper_options.binding_points,
-        std::unordered_map<BindingPoint, core::Access>{}, /* allow_collisions */ true);
+        remapper_data, std::unordered_map<BindingPoint, core::Access>{},
+        /* allow_collisions */ true);
+
+    // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
+    // MultiplanarExternalTexture must come after BindingRemapper
+    data.Add<ast::transform::MultiplanarExternalTexture::NewBindingPoints>(
+        external_texture_options.bindings_map, /* allow_collisions */ true);
+    manager.Add<ast::transform::MultiplanarExternalTexture>();
 
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index a773143..9c9ba7c 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -87,7 +87,7 @@
 };
 
 /// Sanitize a program in preparation for generating MSL.
-/// @program The program to sanitize
+/// @param program The program to sanitize
 /// @param options The MSL generator options.
 /// @returns the sanitized program and any supplementary information
 SanitizedResult Sanitize(const Program& program, const Options& options);
diff --git a/src/tint/lang/msl/writer/ast_printer/helper_test.h b/src/tint/lang/msl/writer/ast_printer/helper_test.h
index 0e9ac5b..b60d232 100644
--- a/src/tint/lang/msl/writer/ast_printer/helper_test.h
+++ b/src/tint/lang/msl/writer/ast_printer/helper_test.h
@@ -106,8 +106,11 @@
   private:
     std::unique_ptr<ASTPrinter> gen_;
 };
+
+/// Test helper
 using TestHelper = TestHelperBase<testing::Test>;
 
+/// Param test helper
 template <typename T>
 using TestParamHelper = TestHelperBase<testing::TestWithParam<T>>;
 
diff --git a/src/tint/lang/msl/writer/common/BUILD.bazel b/src/tint/lang/msl/writer/common/BUILD.bazel
index 3dc8de6..c8f2c0d 100644
--- a/src/tint/lang/msl/writer/common/BUILD.bazel
+++ b/src/tint/lang/msl/writer/common/BUILD.bazel
@@ -39,10 +39,12 @@
 cc_library(
   name = "common",
   srcs = [
+    "option_helpers.cc",
     "options.cc",
     "printer_support.cc",
   ],
   hdrs = [
+    "option_helpers.h",
     "options.h",
     "printer_support.h",
   ],
@@ -52,6 +54,7 @@
     "//src/tint/lang/core",
     "//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",
diff --git a/src/tint/lang/msl/writer/common/BUILD.cmake b/src/tint/lang/msl/writer/common/BUILD.cmake
index 2178ba0..177f390 100644
--- a/src/tint/lang/msl/writer/common/BUILD.cmake
+++ b/src/tint/lang/msl/writer/common/BUILD.cmake
@@ -41,6 +41,8 @@
 # Condition: TINT_BUILD_MSL_WRITER
 ################################################################################
 tint_add_target(tint_lang_msl_writer_common lib
+  lang/msl/writer/common/option_helpers.cc
+  lang/msl/writer/common/option_helpers.h
   lang/msl/writer/common/options.cc
   lang/msl/writer/common/options.h
   lang/msl/writer/common/printer_support.cc
@@ -53,6 +55,7 @@
   tint_lang_core
   tint_lang_core_type
   tint_utils_containers
+  tint_utils_diagnostic
   tint_utils_ice
   tint_utils_id
   tint_utils_macros
diff --git a/src/tint/lang/msl/writer/common/BUILD.gn b/src/tint/lang/msl/writer/common/BUILD.gn
index 287e254..0ea4cc8 100644
--- a/src/tint/lang/msl/writer/common/BUILD.gn
+++ b/src/tint/lang/msl/writer/common/BUILD.gn
@@ -44,6 +44,8 @@
 if (tint_build_msl_writer) {
   libtint_source_set("common") {
     sources = [
+      "option_helpers.cc",
+      "option_helpers.h",
       "options.cc",
       "options.h",
       "printer_support.cc",
@@ -55,6 +57,7 @@
       "${tint_src_dir}/lang/core",
       "${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",
diff --git a/src/tint/lang/spirv/writer/common/option_builder.cc b/src/tint/lang/msl/writer/common/option_helpers.cc
similarity index 60%
copy from src/tint/lang/spirv/writer/common/option_builder.cc
copy to src/tint/lang/msl/writer/common/option_helpers.cc
index 55181fc..cd44e5e 100644
--- a/src/tint/lang/spirv/writer/common/option_builder.cc
+++ b/src/tint/lang/msl/writer/common/option_helpers.cc
@@ -1,4 +1,4 @@
-// Copyright 2023 The Dawn & Tint Authors
+/// 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:
@@ -25,15 +25,21 @@
 // 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/writer/common/option_builder.h"
+#include "src/tint/lang/msl/writer/common/option_helpers.h"
 
 #include "src/tint/utils/containers/hashset.h"
 
-namespace tint::spirv::writer {
+namespace tint::msl::writer {
+
+/// binding::BindingInfo to tint::BindingPoint map
+using InfoToPointMap = tint::Hashmap<binding::BindingInfo, tint::BindingPoint, 8>;
 
 bool ValidateBindingOptions(const Options& options, diag::List& diagnostics) {
     tint::Hashmap<tint::BindingPoint, binding::BindingInfo, 8> seen_wgsl_bindings{};
-    tint::Hashmap<binding::BindingInfo, tint::BindingPoint, 8> seen_spirv_bindings{};
+
+    InfoToPointMap seen_msl_buffer_bindings{};
+    InfoToPointMap seen_msl_texture_bindings{};
+    InfoToPointMap seen_msl_sampler_bindings{};
 
     // Both wgsl_seen and spirv_seen check to see if the pair of [src, dst] are unique. If we have
     // multiple entries that map the same [src, dst] pair, that's fine. We treat it as valid as it's
@@ -55,22 +61,21 @@
         return false;
     };
 
-    auto spirv_seen = [&diagnostics, &seen_spirv_bindings](const binding::BindingInfo& src,
-                                                           const tint::BindingPoint& dst) -> bool {
-        if (auto binding = seen_spirv_bindings.Find(src)) {
+    auto msl_seen = [&diagnostics](InfoToPointMap& map, const binding::BindingInfo& src,
+                                   const tint::BindingPoint& dst) -> bool {
+        if (auto binding = map.Find(src)) {
             if (*binding != dst) {
                 std::stringstream str;
-                str << "found duplicate SPIR-V binding point: [group: " << src.group
-                    << ", binding: " << src.binding << "]";
+                str << "found duplicate MSL binding point: [binding: " << src.binding << "]";
                 diagnostics.add_error(diag::System::Writer, str.str());
                 return true;
             }
         }
-        seen_spirv_bindings.Add(src, dst);
+        map.Add(src, dst);
         return false;
     };
 
-    auto valid = [&wgsl_seen, &spirv_seen](const auto& hsh) -> bool {
+    auto valid = [&wgsl_seen, &msl_seen](InfoToPointMap& map, const auto& hsh) -> bool {
         for (const auto& it : hsh) {
             const auto& src_binding = it.first;
             const auto& dst_binding = it.second;
@@ -79,33 +84,38 @@
                 return false;
             }
 
-            if (spirv_seen(dst_binding, src_binding)) {
+            if (msl_seen(map, dst_binding, src_binding)) {
                 return false;
             }
         }
         return true;
     };
 
-    if (!valid(options.bindings.uniform)) {
+    // Storage and uniform are both [[buffer()]]
+    if (!valid(seen_msl_buffer_bindings, options.bindings.uniform)) {
         diagnostics.add_note(diag::System::Writer, "when processing uniform", {});
         return false;
     }
-    if (!valid(options.bindings.storage)) {
+    if (!valid(seen_msl_buffer_bindings, options.bindings.storage)) {
         diagnostics.add_note(diag::System::Writer, "when processing storage", {});
         return false;
     }
-    if (!valid(options.bindings.texture)) {
+
+    // Sampler is [[sampler()]]
+    if (!valid(seen_msl_sampler_bindings, options.bindings.sampler)) {
+        diagnostics.add_note(diag::System::Writer, "when processing sampler", {});
+        return false;
+    }
+
+    // Texture and storage texture are [[texture()]]
+    if (!valid(seen_msl_texture_bindings, options.bindings.texture)) {
         diagnostics.add_note(diag::System::Writer, "when processing texture", {});
         return false;
     }
-    if (!valid(options.bindings.storage_texture)) {
+    if (!valid(seen_msl_texture_bindings, options.bindings.storage_texture)) {
         diagnostics.add_note(diag::System::Writer, "when processing storage_texture", {});
         return false;
     }
-    if (!valid(options.bindings.sampler)) {
-        diagnostics.add_note(diag::System::Writer, "when processing sampler", {});
-        return false;
-    }
 
     for (const auto& it : options.bindings.external_texture) {
         const auto& src_binding = it.first;
@@ -119,15 +129,17 @@
             return false;
         }
 
-        if (spirv_seen(plane0, src_binding)) {
+        // Plane0 & Plane1 are [[texture()]]
+        if (msl_seen(seen_msl_texture_bindings, plane0, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
-        if (spirv_seen(plane1, src_binding)) {
+        if (msl_seen(seen_msl_texture_bindings, plane1, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
-        if (spirv_seen(metadata, src_binding)) {
+        // Metadata is [[buffer()]]
+        if (msl_seen(seen_msl_buffer_bindings, metadata, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
@@ -139,44 +151,8 @@
 // The remapped binding data and external texture data need to coordinate in order to put things in
 // the correct place when we're done.
 //
-// When the data comes in we have a list of all WGSL origin (group,binding) pairs to SPIR-V
-// (group,binding) pairs in the `uniform`, `storage`, `texture`, and `sampler` arrays.
-//
-// The `external_texture` array stores a WGSL origin (group,binding) pair for the external textures
-// which provide `plane0`, `plane1`, and `metadata` SPIR-V (group,binding) pairs.
-//
-// If the remapper is run first, then the `external_texture` will end up being moved from the WGSL
-// point, or the SPIR-V point (or the `plane0` value). There will also, possibly, have been bindings
-// moved aside in order to place the `external_texture` bindings.
-//
-// If multiplanar runs first, care needs to be taken that when the texture is split and we create
-// `plane1` and `metadata` that they do not collide with existing bindings. If they would collide
-// then we need to place them elsewhere and have the remapper place them in the correct locations.
-//
-// # Example
-// WGSL:
-//   @group(0) @binding(0) var<uniform> u: Uniforms;
-//   @group(0) @binding(1) var s: sampler;
-//   @group(0) @binding(2) var t: texture_external;
-//
-// Given that program, Dawn may decide to do the remappings such that:
-//   * WGSL u (0, 0) -> SPIR-V (0, 1)
-//   * WGSL s (0, 1) -> SPIR-V (0, 2)
-//   * WGSL t (0, 2):
-//     * plane0 -> SPIR-V (0, 3)
-//     * plane1 -> SPIR-V (0, 4)
-//     * metadata -> SPIR-V (0, 0)
-//
-// In this case, if we run binding remapper first, then tell multiplanar to look for the texture at
-// (0, 3) instead of the original (0, 2).
-//
-// If multiplanar runs first, then metadata (0, 0) needs to be placed elsewhere and then remapped
-// back to (0, 0) by the remapper. (Otherwise, we'll have two `@group(0) @binding(0)` items in the
-// program.)
-//
-// # Status
-// The below method assumes we run binding remapper first. So it will setup the binding data and
-// switch the value used by the multiplanar.
+// When the data comes in we have a list of all WGSL origin (group,binding) pairs to MSL
+// (binding) in the `uniform`, `storage`, `texture`, and `sampler` arrays.
 void PopulateRemapperAndMultiplanarOptions(const Options& options,
                                            RemapperData& remapper_data,
                                            ExternalTextureOptions& external_texture) {
@@ -185,14 +161,13 @@
             const BindingPoint& src_binding_point = it.first;
             const binding::Uniform& dst_binding_point = it.second;
 
-            // Bindings which go to the same slot in SPIR-V do not need to be re-bound.
-            if (src_binding_point.group == dst_binding_point.group &&
+            // Bindings which go to the same slot in MSL do not need to be re-bound.
+            if (src_binding_point.group == 0 &&
                 src_binding_point.binding == dst_binding_point.binding) {
                 continue;
             }
 
-            remapper_data.emplace(src_binding_point,
-                                  BindingPoint{dst_binding_point.group, dst_binding_point.binding});
+            remapper_data.emplace(src_binding_point, BindingPoint{0, dst_binding_point.binding});
         }
     };
 
@@ -205,42 +180,42 @@
     // External textures are re-bound to their plane0 location
     for (const auto& it : options.bindings.external_texture) {
         const BindingPoint& src_binding_point = it.first;
+
         const binding::BindingInfo& plane0 = it.second.plane0;
         const binding::BindingInfo& plane1 = it.second.plane1;
         const binding::BindingInfo& metadata = it.second.metadata;
 
-        BindingPoint plane0_binding_point{plane0.group, plane0.binding};
-        BindingPoint plane1_binding_point{plane1.group, plane1.binding};
-        BindingPoint metadata_binding_point{metadata.group, metadata.binding};
+        BindingPoint plane0_binding_point{0, plane0.binding};
+        BindingPoint plane1_binding_point{0, plane1.binding};
+        BindingPoint metadata_binding_point{0, metadata.binding};
 
-        // Use the re-bound spir-v plane0 value for the lookup key.
+        // Use the re-bound msl plane0 value for the lookup key.
         external_texture.bindings_map.emplace(
-            plane0_binding_point,
+            BindingPoint{src_binding_point.group, plane0_binding_point.binding},
             ExternalTextureOptions::BindingPoints{plane1_binding_point, metadata_binding_point});
 
-        // Bindings which go to the same slot in SPIR-V do not need to be re-bound.
-        if (src_binding_point.group == plane0.group &&
-            src_binding_point.binding == plane0.binding) {
+        // Bindings which go to the same slot in MSL do not need to be re-bound.
+        if (src_binding_point == plane0_binding_point) {
             continue;
         }
 
-        remapper_data.emplace(src_binding_point, BindingPoint{plane0.group, plane0.binding});
+        remapper_data.emplace(src_binding_point, plane0_binding_point);
     }
 }
 
-}  // namespace tint::spirv::writer
+}  // namespace tint::msl::writer
 
 namespace std {
 
-/// Custom std::hash specialization for tint::spirv::writer::binding::BindingInfo so
+/// Custom std::hash specialization for tint::msl::writer::binding::BindingInfo so
 /// they can be used as keys for std::unordered_map and std::unordered_set.
 template <>
-class hash<tint::spirv::writer::binding::BindingInfo> {
+class hash<tint::msl::writer::binding::BindingInfo> {
   public:
     /// @param info the binding to create a hash for
     /// @return the hash value
-    inline std::size_t operator()(const tint::spirv::writer::binding::BindingInfo& info) const {
-        return tint::Hash(info.group, info.binding);
+    inline std::size_t operator()(const tint::msl::writer::binding::BindingInfo& info) const {
+        return tint::Hash(info.binding);
     }
 };
 
diff --git a/src/tint/lang/spirv/writer/common/option_builder.h b/src/tint/lang/msl/writer/common/option_helpers.h
similarity index 85%
copy from src/tint/lang/spirv/writer/common/option_builder.h
copy to src/tint/lang/msl/writer/common/option_helpers.h
index decedd3..1abf984 100644
--- a/src/tint/lang/spirv/writer/common/option_builder.h
+++ b/src/tint/lang/msl/writer/common/option_helpers.h
@@ -1,4 +1,4 @@
-// Copyright 2023 The Dawn & Tint Authors
+/// 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:
@@ -25,21 +25,23 @@
 // 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_WRITER_COMMON_OPTION_BUILDER_H_
-#define SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_BUILDER_H_
+#ifndef SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_HELPERS_H_
+#define SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_HELPERS_H_
 
 #include <unordered_map>
 
 #include "src/tint/api/common/binding_point.h"
 #include "src/tint/api/options/external_texture.h"
-#include "src/tint/lang/spirv/writer/common/options.h"
+#include "src/tint/lang/msl/writer/common/options.h"
 #include "src/tint/utils/diagnostic/diagnostic.h"
 
-namespace tint::spirv::writer {
+namespace tint::msl::writer {
 
+/// The remapper data
 using RemapperData = std::unordered_map<BindingPoint, BindingPoint>;
 
 /// @param options the options
+/// @param diagnostics the diagnostics
 /// @returns true if the binding points are valid
 bool ValidateBindingOptions(const Options& options, diag::List& diagnostics);
 
@@ -52,6 +54,6 @@
                                            RemapperData& remapper_data,
                                            ExternalTextureOptions& external_texture);
 
-}  // namespace tint::spirv::writer
+}  // namespace tint::msl::writer
 
-#endif  // SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_BUILDER_H_
+#endif  // SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_HELPERS_H_
diff --git a/src/tint/lang/msl/writer/common/options.h b/src/tint/lang/msl/writer/common/options.h
index 6ce31f1..d085a66 100644
--- a/src/tint/lang/msl/writer/common/options.h
+++ b/src/tint/lang/msl/writer/common/options.h
@@ -28,13 +28,85 @@
 #ifndef SRC_TINT_LANG_MSL_WRITER_COMMON_OPTIONS_H_
 #define SRC_TINT_LANG_MSL_WRITER_COMMON_OPTIONS_H_
 
+#include <unordered_map>
+
+#include "src/tint/api/common/binding_point.h"
 #include "src/tint/api/options/array_length_from_uniform.h"
-#include "src/tint/api/options/binding_remapper.h"
-#include "src/tint/api/options/external_texture.h"
 #include "src/tint/api/options/pixel_local.h"
 #include "src/tint/utils/reflection/reflection.h"
 
 namespace tint::msl::writer {
+namespace binding {
+
+/// Generic binding point
+struct BindingInfo {
+    /// The binding
+    uint32_t binding = 0;
+
+    /// Equality operator
+    /// @param rhs the BindingInfo to compare against
+    /// /// @returns true if this BindingInfo is equal to `rhs`
+    inline bool operator==(const BindingInfo& rhs) const { return binding == rhs.binding; }
+    /// Inequality operator
+    /// @param rhs the BindingInfo to compare against
+    /// @returns true if this BindingInfo is not equal to `rhs`
+    inline bool operator!=(const BindingInfo& rhs) const { return !(*this == rhs); }
+
+    /// Reflect the fields of this class so taht it can be used by tint::ForeachField()
+    TINT_REFLECT(binding);
+};
+using Uniform = BindingInfo;
+using Storage = BindingInfo;
+using Texture = BindingInfo;
+using StorageTexture = BindingInfo;
+using Sampler = BindingInfo;
+
+/// An external texture
+struct ExternalTexture {
+    /// Metadata
+    BindingInfo metadata{};
+    /// Plane0 binding data
+    BindingInfo plane0{};
+    /// Plane1 binding data;
+    BindingInfo plane1{};
+
+    /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+    TINT_REFLECT(metadata, plane0, plane1);
+};
+
+}  // namespace binding
+
+/// Maps the WGSL binding point to the SPIR-V group,binding for uniforms
+using UniformBindings = std::unordered_map<BindingPoint, binding::Uniform>;
+/// Maps the WGSL binding point to the SPIR-V group,binding for storage
+using StorageBindings = std::unordered_map<BindingPoint, binding::Storage>;
+/// Maps the WGSL binding point to the SPIR-V group,binding for textures
+using TextureBindings = std::unordered_map<BindingPoint, binding::Texture>;
+/// Maps the WGSL binding point to the SPIR-V group,binding for storage textures
+using StorageTextureBindings = std::unordered_map<BindingPoint, binding::StorageTexture>;
+/// Maps the WGSL binding point to the SPIR-V group,binding for samplers
+using SamplerBindings = std::unordered_map<BindingPoint, binding::Sampler>;
+/// Maps the WGSL binding point to the plane0, plane1, and metadata for external textures
+using ExternalTextureBindings = std::unordered_map<BindingPoint, binding::ExternalTexture>;
+
+/// Binding information
+struct Bindings {
+    /// Uniform bindings
+    UniformBindings uniform{};
+    /// Storage bindings
+    StorageBindings storage{};
+    /// Texture bindings
+    TextureBindings texture{};
+    /// Storage texture bindings
+    StorageTextureBindings storage_texture{};
+    /// Sampler bindings
+    SamplerBindings sampler{};
+    /// External bindings
+    ExternalTextureBindings external_texture{};
+
+    /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+    TINT_REFLECT(uniform, storage, texture, storage_texture, sampler, external_texture);
+};
 
 /// Configuration options used for generating MSL.
 struct Options {
@@ -76,11 +148,8 @@
     /// from which to load buffer sizes.
     ArrayLengthFromUniformOptions array_length_from_uniform = {};
 
-    /// Options used in the binding mappings for external textures
-    ExternalTextureOptions external_texture_options = {};
-
-    /// Options used in the bindings remapper
-    BindingRemapperOptions binding_remapper_options = {};
+    /// The bindings
+    Bindings bindings;
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
     TINT_REFLECT(disable_robustness,
@@ -91,8 +160,7 @@
                  fixed_sample_mask,
                  pixel_local_options,
                  array_length_from_uniform,
-                 external_texture_options,
-                 binding_remapper_options);
+                 bindings);
 };
 
 }  // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/helpers/BUILD.bazel b/src/tint/lang/msl/writer/helpers/BUILD.bazel
new file mode 100644
index 0000000..095cd37
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/BUILD.bazel
@@ -0,0 +1,84 @@
+# 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 = "helpers",
+  srcs = [
+    "generate_bindings.cc",
+  ],
+  hdrs = [
+    "generate_bindings.h",
+  ],
+  deps = [
+    "//src/tint/api/common",
+    "//src/tint/api/options",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/type",
+    "//src/tint/lang/wgsl",
+    "//src/tint/lang/wgsl/ast",
+    "//src/tint/lang/wgsl/program",
+    "//src/tint/lang/wgsl/sem",
+    "//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({
+    ":tint_build_msl_writer": [
+      "//src/tint/lang/msl/writer/common",
+    ],
+    "//conditions:default": [],
+  }),
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
+
+alias(
+  name = "tint_build_msl_writer",
+  actual = "//src/tint:tint_build_msl_writer_true",
+)
+
diff --git a/src/tint/lang/msl/writer/helpers/BUILD.cfg b/src/tint/lang/msl/writer/helpers/BUILD.cfg
new file mode 100644
index 0000000..1c7e256
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/BUILD.cfg
@@ -0,0 +1,3 @@
+{
+    "condition": "tint_build_msl_writer"
+}
diff --git a/src/tint/lang/msl/writer/helpers/BUILD.cmake b/src/tint/lang/msl/writer/helpers/BUILD.cmake
new file mode 100644
index 0000000..1e27fed
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/BUILD.cmake
@@ -0,0 +1,79 @@
+# 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
+################################################################################
+
+if(TINT_BUILD_MSL_WRITER)
+################################################################################
+# Target:    tint_lang_msl_writer_helpers
+# Kind:      lib
+# Condition: TINT_BUILD_MSL_WRITER
+################################################################################
+tint_add_target(tint_lang_msl_writer_helpers lib
+  lang/msl/writer/helpers/generate_bindings.cc
+  lang/msl/writer/helpers/generate_bindings.h
+)
+
+tint_target_add_dependencies(tint_lang_msl_writer_helpers lib
+  tint_api_common
+  tint_api_options
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_type
+  tint_lang_wgsl
+  tint_lang_wgsl_ast
+  tint_lang_wgsl_program
+  tint_lang_wgsl_sem
+  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
+)
+
+if(TINT_BUILD_MSL_WRITER)
+  tint_target_add_dependencies(tint_lang_msl_writer_helpers lib
+    tint_lang_msl_writer_common
+  )
+endif(TINT_BUILD_MSL_WRITER)
+
+endif(TINT_BUILD_MSL_WRITER)
\ No newline at end of file
diff --git a/src/tint/lang/msl/writer/helpers/BUILD.gn b/src/tint/lang/msl/writer/helpers/BUILD.gn
new file mode 100644
index 0000000..3b0713b
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/BUILD.gn
@@ -0,0 +1,75 @@
+# 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_msl_writer) {
+  libtint_source_set("helpers") {
+    sources = [
+      "generate_bindings.cc",
+      "generate_bindings.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/type",
+      "${tint_src_dir}/lang/wgsl",
+      "${tint_src_dir}/lang/wgsl/ast",
+      "${tint_src_dir}/lang/wgsl/program",
+      "${tint_src_dir}/lang/wgsl/sem",
+      "${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/common" ]
+    }
+  }
+}
diff --git a/src/tint/lang/msl/writer/helpers/generate_bindings.cc b/src/tint/lang/msl/writer/helpers/generate_bindings.cc
new file mode 100644
index 0000000..a0eaf7e
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/generate_bindings.cc
@@ -0,0 +1,126 @@
+/// 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/helpers/generate_bindings.h"
+
+#include <algorithm>
+#include <unordered_map>
+#include <unordered_set>
+#include <vector>
+
+#include "src/tint/api/common/binding_point.h"
+#include "src/tint/lang/core/type/external_texture.h"
+#include "src/tint/lang/core/type/pointer.h"
+#include "src/tint/lang/core/type/storage_texture.h"
+#include "src/tint/lang/msl/writer/common/options.h"
+#include "src/tint/lang/wgsl/ast/module.h"
+#include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/lang/wgsl/sem/variable.h"
+#include "src/tint/utils/ice/ice.h"
+#include "src/tint/utils/rtti/switch.h"
+
+namespace tint::msl::writer {
+
+Bindings GenerateBindings(const Program& program) {
+    // TODO(tint:1491): Use Inspector once we can get binding info for all
+    // variables, not just those referenced by entry points.
+
+    Bindings bindings{};
+
+    std::unordered_set<tint::BindingPoint> seen_binding_points;
+
+    // Collect next valid binding number per group
+    Hashmap<uint32_t, uint32_t, 4> group_to_next_binding_number;
+    Vector<tint::BindingPoint, 4> ext_tex_bps;
+    for (auto* var : program.AST().GlobalVariables()) {
+        if (auto* sem_var = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
+            if (auto bp = sem_var->Attributes().binding_point) {
+                if (auto val = group_to_next_binding_number.Find(bp->group)) {
+                    *val = std::max(*val, bp->binding + 1);
+                } else {
+                    group_to_next_binding_number.Add(bp->group, bp->binding + 1);
+                }
+
+                // Store up the external textures, we'll add them in the next step
+                if (sem_var->Type()->UnwrapRef()->Is<core::type::ExternalTexture>()) {
+                    ext_tex_bps.Push(*bp);
+                    continue;
+                }
+
+                binding::BindingInfo info{bp->binding};
+                switch (sem_var->AddressSpace()) {
+                    case core::AddressSpace::kHandle:
+                        Switch(
+                            sem_var->Type()->UnwrapRef(),  //
+                            [&](const core::type::Sampler*) {
+                                bindings.sampler.emplace(*bp, info);
+                            },
+                            [&](const core::type::StorageTexture*) {
+                                bindings.storage_texture.emplace(*bp, info);
+                            },
+                            [&](const core::type::Texture*) {
+                                bindings.texture.emplace(*bp, info);
+                            });
+                        break;
+                    case core::AddressSpace::kStorage:
+                        bindings.storage.emplace(*bp, info);
+                        break;
+                    case core::AddressSpace::kUniform:
+                        bindings.uniform.emplace(*bp, info);
+                        break;
+
+                    case core::AddressSpace::kUndefined:
+                    case core::AddressSpace::kPixelLocal:
+                    case core::AddressSpace::kPrivate:
+                    case core::AddressSpace::kPushConstant:
+                    case core::AddressSpace::kIn:
+                    case core::AddressSpace::kOut:
+                    case core::AddressSpace::kFunction:
+                    case core::AddressSpace::kWorkgroup:
+                        break;
+                }
+            }
+        }
+    }
+
+    for (auto bp : ext_tex_bps) {
+        uint32_t g = bp.group;
+        uint32_t next_num = *(group_to_next_binding_number.GetOrZero(g));
+
+        binding::BindingInfo plane0{bp.binding};
+        binding::BindingInfo plane1{next_num++};
+        binding::BindingInfo metadata{next_num++};
+
+        group_to_next_binding_number.Replace(g, next_num);
+
+        bindings.external_texture.emplace(bp, binding::ExternalTexture{metadata, plane0, plane1});
+    }
+
+    return bindings;
+}
+
+}  // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/helpers/generate_bindings.h b/src/tint/lang/msl/writer/helpers/generate_bindings.h
new file mode 100644
index 0000000..9cfc85f
--- /dev/null
+++ b/src/tint/lang/msl/writer/helpers/generate_bindings.h
@@ -0,0 +1,47 @@
+/// 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_HELPERS_GENERATE_BINDINGS_H_
+#define SRC_TINT_LANG_MSL_WRITER_HELPERS_GENERATE_BINDINGS_H_
+
+#include "src/tint/lang/msl/writer/common/options.h"
+
+// Forward declarations
+namespace tint {
+class Program;
+}
+
+namespace tint::msl::writer {
+
+/// Generate the resource bindings
+/// @param program the program to generate from
+/// @returns the bindings
+Bindings GenerateBindings(const Program& program);
+
+}  // namespace tint::msl::writer
+
+#endif  // SRC_TINT_LANG_MSL_WRITER_HELPERS_GENERATE_BINDINGS_H_
diff --git a/src/tint/lang/msl/writer/printer/helper_test.h b/src/tint/lang/msl/writer/printer/helper_test.h
index e0ecbf8..f77774b 100644
--- a/src/tint/lang/msl/writer/printer/helper_test.h
+++ b/src/tint/lang/msl/writer/printer/helper_test.h
@@ -39,10 +39,12 @@
 
 namespace tint::msl::writer {
 
+/// Metal header declaration
 constexpr auto kMetalHeader = R"(#include <metal_stdlib>
 using namespace metal;
 )";
 
+/// Metal array declaration
 constexpr auto kMetalArray = R"(template<typename T, size_t N>
 struct tint_array {
   const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -99,8 +101,10 @@
     std::string MetalArray() const { return kMetalArray; }
 };
 
+/// Printer tests
 using MslPrinterTest = MslPrinterTestHelperBase<testing::Test>;
 
+/// Printer param tests
 template <typename T>
 using MslPrinterTestWithParam = MslPrinterTestHelperBase<testing::TestWithParam<T>>;
 
diff --git a/src/tint/lang/msl/writer/writer.cc b/src/tint/lang/msl/writer/writer.cc
index 78f3040..cb5281d 100644
--- a/src/tint/lang/msl/writer/writer.cc
+++ b/src/tint/lang/msl/writer/writer.cc
@@ -31,6 +31,7 @@
 #include <utility>
 
 #include "src/tint/lang/msl/writer/ast_printer/ast_printer.h"
+#include "src/tint/lang/msl/writer/common/option_helpers.h"
 #include "src/tint/lang/msl/writer/printer/printer.h"
 #include "src/tint/lang/msl/writer/raise/raise.h"
 
@@ -46,6 +47,13 @@
         return Failure{program.Diagnostics()};
     }
 
+    {
+        diag::List validation_diagnostics;
+        if (!ValidateBindingOptions(options, validation_diagnostics)) {
+            return Failure{validation_diagnostics};
+        }
+    }
+
     Output output;
 
     if (options.use_tint_ir) {
diff --git a/src/tint/lang/msl/writer/writer_bench.cc b/src/tint/lang/msl/writer/writer_bench.cc
index 0af1f7b..b89b8f7 100644
--- a/src/tint/lang/msl/writer/writer_bench.cc
+++ b/src/tint/lang/msl/writer/writer_bench.cc
@@ -28,6 +28,7 @@
 #include <string>
 
 #include "src/tint/cmd/bench/bench.h"
+#include "src/tint/lang/msl/writer/helpers/generate_bindings.h"
 #include "src/tint/lang/msl/writer/writer.h"
 #include "src/tint/lang/wgsl/ast/module.h"
 #include "src/tint/lang/wgsl/sem/variable.h"
@@ -61,18 +62,8 @@
                                                                           6);
     gen_options.array_length_from_uniform.bindpoint_to_size_index.emplace(tint::BindingPoint{0, 7},
                                                                           7);
+    gen_options.bindings = tint::msl::writer::GenerateBindings(res->program);
 
-    uint32_t next_binding_point = 0;
-    for (auto* var : program.AST().GlobalVariables()) {
-        if (auto* var_sem = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
-            if (auto bp = var_sem->Attributes().binding_point) {
-                gen_options.binding_remapper_options.binding_points[*bp] = BindingPoint{
-                    0,                     // group
-                    next_binding_point++,  // binding
-                };
-            }
-        }
-    }
     for (auto _ : state) {
         auto gen_res = Generate(program, gen_options);
         if (!gen_res) {
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
index 9afa797..61ceab1 100644
--- a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
@@ -37,7 +37,7 @@
 #include "src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index.h"
 #include "src/tint/lang/spirv/writer/ast_raise/vectorize_matrix_conversions.h"
 #include "src/tint/lang/spirv/writer/ast_raise/while_to_loop.h"
-#include "src/tint/lang/spirv/writer/common/option_builder.h"
+#include "src/tint/lang/spirv/writer/common/option_helpers.h"
 #include "src/tint/lang/wgsl/ast/transform/add_block_attribute.h"
 #include "src/tint/lang/wgsl/ast/transform/add_empty_entry_point.h"
 #include "src/tint/lang/wgsl/ast/transform/binding_remapper.h"
diff --git a/src/tint/lang/spirv/writer/common/BUILD.bazel b/src/tint/lang/spirv/writer/common/BUILD.bazel
index 7a3f041..7c07fad 100644
--- a/src/tint/lang/spirv/writer/common/BUILD.bazel
+++ b/src/tint/lang/spirv/writer/common/BUILD.bazel
@@ -44,7 +44,7 @@
     "instruction.cc",
     "module.cc",
     "operand.cc",
-    "option_builder.cc",
+    "option_helper.cc",
   ],
   hdrs = [
     "binary_writer.h",
@@ -52,7 +52,7 @@
     "instruction.h",
     "module.h",
     "operand.h",
-    "option_builder.h",
+    "option_helpers.h",
     "options.h",
   ],
   deps = [
diff --git a/src/tint/lang/spirv/writer/common/BUILD.cmake b/src/tint/lang/spirv/writer/common/BUILD.cmake
index 94f5a1b..969c5e7 100644
--- a/src/tint/lang/spirv/writer/common/BUILD.cmake
+++ b/src/tint/lang/spirv/writer/common/BUILD.cmake
@@ -51,8 +51,8 @@
   lang/spirv/writer/common/module.h
   lang/spirv/writer/common/operand.cc
   lang/spirv/writer/common/operand.h
-  lang/spirv/writer/common/option_builder.cc
-  lang/spirv/writer/common/option_builder.h
+  lang/spirv/writer/common/option_helper.cc
+  lang/spirv/writer/common/option_helpers.h
   lang/spirv/writer/common/options.h
 )
 
diff --git a/src/tint/lang/spirv/writer/common/BUILD.gn b/src/tint/lang/spirv/writer/common/BUILD.gn
index 53c6f01..0c9830d 100644
--- a/src/tint/lang/spirv/writer/common/BUILD.gn
+++ b/src/tint/lang/spirv/writer/common/BUILD.gn
@@ -54,8 +54,8 @@
       "module.h",
       "operand.cc",
       "operand.h",
-      "option_builder.cc",
-      "option_builder.h",
+      "option_helper.cc",
+      "option_helpers.h",
       "options.h",
     ]
     deps = [
diff --git a/src/tint/lang/spirv/writer/common/option_builder.cc b/src/tint/lang/spirv/writer/common/option_helper.cc
similarity index 99%
rename from src/tint/lang/spirv/writer/common/option_builder.cc
rename to src/tint/lang/spirv/writer/common/option_helper.cc
index 55181fc..dc3670b 100644
--- a/src/tint/lang/spirv/writer/common/option_builder.cc
+++ b/src/tint/lang/spirv/writer/common/option_helper.cc
@@ -25,7 +25,7 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/spirv/writer/common/option_builder.h"
+#include "src/tint/lang/spirv/writer/common/option_helpers.h"
 
 #include "src/tint/utils/containers/hashset.h"
 
diff --git a/src/tint/lang/spirv/writer/common/option_builder.h b/src/tint/lang/spirv/writer/common/option_helpers.h
similarity index 93%
rename from src/tint/lang/spirv/writer/common/option_builder.h
rename to src/tint/lang/spirv/writer/common/option_helpers.h
index decedd3..0be6d35 100644
--- a/src/tint/lang/spirv/writer/common/option_builder.h
+++ b/src/tint/lang/spirv/writer/common/option_helpers.h
@@ -25,8 +25,8 @@
 // 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_WRITER_COMMON_OPTION_BUILDER_H_
-#define SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_BUILDER_H_
+#ifndef SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_HELPERS_H_
+#define SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_HELPERS_H_
 
 #include <unordered_map>
 
@@ -54,4 +54,4 @@
 
 }  // namespace tint::spirv::writer
 
-#endif  // SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_BUILDER_H_
+#endif  // SRC_TINT_LANG_SPIRV_WRITER_COMMON_OPTION_HELPERS_H_
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h
index 96b585e..46718d2 100644
--- a/src/tint/lang/spirv/writer/common/options.h
+++ b/src/tint/lang/spirv/writer/common/options.h
@@ -107,7 +107,7 @@
     ExternalTextureBindings external_texture{};
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
-    TINT_REFLECT(uniform, storage, texture, sampler, external_texture);
+    TINT_REFLECT(uniform, storage, texture, storage_texture, sampler, external_texture);
 };
 
 /// Configuration options used for generating SPIR-V.
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc
index a5b0bc4..8160fe6 100644
--- a/src/tint/lang/spirv/writer/raise/raise.cc
+++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -45,7 +45,7 @@
 #include "src/tint/lang/core/ir/transform/std140.h"
 #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/spirv/writer/common/option_builder.h"
+#include "src/tint/lang/spirv/writer/common/option_helpers.h"
 #include "src/tint/lang/spirv/writer/raise/builtin_polyfill.h"
 #include "src/tint/lang/spirv/writer/raise/expand_implicit_splats.h"
 #include "src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.h"
diff --git a/src/tint/lang/spirv/writer/writer.cc b/src/tint/lang/spirv/writer/writer.cc
index 1c44959..e6f3098 100644
--- a/src/tint/lang/spirv/writer/writer.cc
+++ b/src/tint/lang/spirv/writer/writer.cc
@@ -31,7 +31,7 @@
 #include <utility>
 
 #include "src/tint/lang/spirv/writer/ast_printer/ast_printer.h"
-#include "src/tint/lang/spirv/writer/common/option_builder.h"
+#include "src/tint/lang/spirv/writer/common/option_helpers.h"
 #include "src/tint/lang/spirv/writer/printer/printer.h"
 #include "src/tint/lang/spirv/writer/raise/raise.h"
 #include "src/tint/lang/wgsl/reader/lower/lower.h"
diff --git a/src/tint/lang/wgsl/ast/transform/binding_remapper.cc b/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
index b2ec92a..842e9ce 100644
--- a/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
+++ b/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
@@ -75,39 +75,6 @@
         return SkipTransform;
     }
 
-    // A set of post-remapped binding points that need to be decorated with a
-    // DisableValidationAttribute to disable binding-point-collision validation
-    std::unordered_set<BindingPoint> add_collision_attr;
-
-    if (remappings->allow_collisions) {
-        // Scan for binding point collisions generated by this transform.
-        // Populate all collisions in the `add_collision_attr` set.
-        for (auto* func_ast : src.AST().Functions()) {
-            if (!func_ast->IsEntryPoint()) {
-                continue;
-            }
-            auto* func = src.Sem().Get(func_ast);
-            std::unordered_map<BindingPoint, int> binding_point_counts;
-            for (auto* global : func->TransitivelyReferencedGlobals()) {
-                if (auto from = global->Attributes().binding_point) {
-                    auto bp_it = remappings->binding_points.find(*from);
-                    if (bp_it != remappings->binding_points.end()) {
-                        // Remapped
-                        BindingPoint to = bp_it->second;
-                        if (binding_point_counts[to]++) {
-                            add_collision_attr.emplace(to);
-                        }
-                    } else {
-                        // No remapping
-                        if (binding_point_counts[*from]++) {
-                            add_collision_attr.emplace(*from);
-                        }
-                    }
-                }
-            }
-        }
-    }
-
     for (auto* var : src.AST().Globals<Var>()) {
         if (var->HasBindingPoint()) {
             auto* global_sem = src.Sem().Get<sem::GlobalVariable>(var);
@@ -134,6 +101,11 @@
                 ctx.Replace(old_binding, new_binding);
                 bp = to;
             }
+            // Add `DisableValidationAttribute`s if required
+            if (remappings->allow_collisions) {
+                auto* attribute = b.Disable(DisabledValidation::kBindingPointCollision);
+                ctx.InsertBefore(var->attributes, *var->attributes.begin(), attribute);
+            }
 
             // Replace any access controls.
             auto ac_it = remappings->access_controls.find(from);
@@ -166,12 +138,6 @@
                                   ctx.Clone(var->attributes));             // attributes
                 ctx.Replace(var, new_var);
             }
-
-            // Add `DisableValidationAttribute`s if required
-            if (add_collision_attr.count(bp)) {
-                auto* attribute = b.Disable(DisabledValidation::kBindingPointCollision);
-                ctx.InsertBefore(var->attributes, *var->attributes.begin(), attribute);
-            }
         }
     }
 
diff --git a/src/tint/lang/wgsl/ast/transform/binding_remapper_test.cc b/src/tint/lang/wgsl/ast/transform/binding_remapper_test.cc
index a3c1408..1fb96a2 100644
--- a/src/tint/lang/wgsl/ast/transform/binding_remapper_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/binding_remapper_test.cc
@@ -116,9 +116,9 @@
   a : f32,
 }
 
-@group(1) @binding(2) var<storage, read> a : S;
+@internal(disable_validation__binding_point_collision) @group(1) @binding(2) var<storage, read> a : S;
 
-@group(3) @binding(2) var<storage, read> b : S;
+@internal(disable_validation__binding_point_collision) @group(3) @binding(2) var<storage, read> b : S;
 
 @compute @workgroup_size(1)
 fn f() {
@@ -160,11 +160,11 @@
   a : f32,
 }
 
-@group(2) @binding(1) var<storage, read_write> a : S;
+@internal(disable_validation__binding_point_collision) @group(2) @binding(1) var<storage, read_write> a : S;
 
-@group(3) @binding(2) var<storage, read_write> b : S;
+@internal(disable_validation__binding_point_collision) @group(3) @binding(2) var<storage, read_write> b : S;
 
-@group(4) @binding(3) var<storage, read> c : S;
+@internal(disable_validation__binding_point_collision) @group(4) @binding(3) var<storage, read> c : S;
 
 @compute @workgroup_size(1)
 fn f() {
@@ -204,9 +204,9 @@
   a : f32,
 }
 
-@group(4) @binding(5) var<storage, read_write> a : S;
+@internal(disable_validation__binding_point_collision) @group(4) @binding(5) var<storage, read_write> a : S;
 
-@group(6) @binding(7) var<storage, read_write> b : S;
+@internal(disable_validation__binding_point_collision) @group(6) @binding(7) var<storage, read_write> b : S;
 
 @compute @workgroup_size(1)
 fn f() {
@@ -310,13 +310,13 @@
   i : i32,
 }
 
-@group(1) @binding(1) var<storage, read> a : S;
+@internal(disable_validation__binding_point_collision) @group(1) @binding(1) var<storage, read> a : S;
 
-@group(1) @binding(1) var<storage, read> b : S;
+@internal(disable_validation__binding_point_collision) @group(1) @binding(1) var<storage, read> b : S;
 
-@group(5) @binding(4) var<storage, read> c : S;
+@internal(disable_validation__binding_point_collision) @group(5) @binding(4) var<storage, read> c : S;
 
-@group(5) @binding(4) var<storage, read> d : S;
+@internal(disable_validation__binding_point_collision) @group(5) @binding(4) var<storage, read> d : S;
 
 @compute @workgroup_size(1)
 fn f1() {
diff --git a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
index 0a5f85e..bdcc60b 100644
--- a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
+++ b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
@@ -151,12 +151,27 @@
             auto& syms = new_binding_symbols[sem_var];
             syms.plane_0 = ctx.Clone(global->name->symbol);
             syms.plane_1 = b.Symbols().New("ext_tex_plane_1");
-            b.GlobalVar(syms.plane_1,
-                        b.ty.sampled_texture(core::type::TextureDimension::k2d, b.ty.f32()),
-                        b.Group(AInt(bps.plane_1.group)), b.Binding(AInt(bps.plane_1.binding)));
+            if (new_binding_points->allow_collisions) {
+                b.GlobalVar(syms.plane_1,
+                            b.ty.sampled_texture(core::type::TextureDimension::k2d, b.ty.f32()),
+                            b.Disable(DisabledValidation::kBindingPointCollision),
+                            b.Group(AInt(bps.plane_1.group)), b.Binding(AInt(bps.plane_1.binding)));
+            } else {
+                b.GlobalVar(syms.plane_1,
+                            b.ty.sampled_texture(core::type::TextureDimension::k2d, b.ty.f32()),
+                            b.Group(AInt(bps.plane_1.group)), b.Binding(AInt(bps.plane_1.binding)));
+            }
             syms.params = b.Symbols().New("ext_tex_params");
-            b.GlobalVar(syms.params, b.ty("ExternalTextureParams"), core::AddressSpace::kUniform,
-                        b.Group(AInt(bps.params.group)), b.Binding(AInt(bps.params.binding)));
+            if (new_binding_points->allow_collisions) {
+                b.GlobalVar(syms.params, b.ty("ExternalTextureParams"),
+                            core::AddressSpace::kUniform,
+                            b.Disable(DisabledValidation::kBindingPointCollision),
+                            b.Group(AInt(bps.params.group)), b.Binding(AInt(bps.params.binding)));
+            } else {
+                b.GlobalVar(syms.params, b.ty("ExternalTextureParams"),
+                            core::AddressSpace::kUniform, b.Group(AInt(bps.params.group)),
+                            b.Binding(AInt(bps.params.binding)));
+            }
 
             // Replace the original texture_external binding with a texture_2d<f32> binding.
             auto cloned_attributes = ctx.Clone(global->attributes);
@@ -511,8 +526,9 @@
     }
 };
 
-MultiplanarExternalTexture::NewBindingPoints::NewBindingPoints(BindingsMap inputBindingsMap)
-    : bindings_map(std::move(inputBindingsMap)) {}
+MultiplanarExternalTexture::NewBindingPoints::NewBindingPoints(BindingsMap inputBindingsMap,
+                                                               bool may_collide)
+    : bindings_map(std::move(inputBindingsMap)), allow_collisions(may_collide) {}
 
 MultiplanarExternalTexture::NewBindingPoints::~NewBindingPoints() = default;
 
diff --git a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.h b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.h
index 8770b6c..3395047 100644
--- a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.h
+++ b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.h
@@ -67,13 +67,19 @@
     struct NewBindingPoints final : public Castable<NewBindingPoints, Data> {
         /// Constructor
         /// @param bm a map to the new binding slots to use.
-        explicit NewBindingPoints(BindingsMap bm);
+        /// /// @param may_collide if true, then validation will be disabled for binding point
+        /// collisions generated by this transform
+        explicit NewBindingPoints(BindingsMap bm, bool may_collide = false);
 
         /// Destructor
         ~NewBindingPoints() override;
 
         /// A map of new binding points to use.
         const BindingsMap bindings_map;
+
+        /// If true, then validation will be disabled for bindign poitn collisions generated by this
+        /// transform.
+        const bool allow_collisions = false;
     };
 
     /// Constructor