Import Tint changes from Dawn
Changes:
- 35dbf7f4cc7f0cd89231ce5f26d38e20b00e9aea Rename option_builder to option_helpers. by dan sinclair <dsinclair@chromium.org>
- 4c88c1eade1b75def67979a37b152e8329c3c34b Using binding information for MSL/Tint interface by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: 35dbf7f4cc7f0cd89231ce5f26d38e20b00e9aea
Change-Id: Ifdd8964520860a5290d59356837444e389be1051
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/160683
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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
