Using binding information for MSL/Tint interface
This CL removes the explicit external texture and multiplanar options
from the MSL writer and instead uses `binding` information which is
populated by Dawn.
On the Tint side the needed transform data is created from the requested
result from Dawn.
Change-Id: I56eeb33cb09435e015a7c9e5a6021a20726b1c2a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/156761
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index a1553fb..2a770fc 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -27,10 +27,11 @@
#include "dawn/native/metal/ShaderModuleMTL.h"
-#include "dawn/native/BindGroupLayoutInternal.h"
+#include "dawn/native/BindGroupLayout.h"
#include "dawn/native/CacheRequest.h"
#include "dawn/native/Serializable.h"
#include "dawn/native/TintUtils.h"
+#include "dawn/native/metal/BindGroupLayoutMTL.h"
#include "dawn/native/metal/DeviceMTL.h"
#include "dawn/native/metal/PipelineLayoutMTL.h"
#include "dawn/native/metal/RenderPipelineMTL.h"
@@ -114,39 +115,30 @@
const PipelineLayout* layout,
ShaderModule::MetalFunctionData* out,
uint32_t sampleMask,
- const RenderPipeline* renderPipeline) {
+ const RenderPipeline* renderPipeline,
+ const BindingInfoArray& moduleBindingInfo) {
ScopedTintICEHandler scopedICEHandler(device);
std::ostringstream errorStream;
errorStream << "Tint MSL failure:" << std::endl;
- // Remap BindingNumber to BindingIndex in WGSL shader
- using BindingPoint = tint::BindingPoint;
-
- tint::BindingRemapperOptions bindingRemapper;
-
tint::ArrayLengthFromUniformOptions arrayLengthFromUniform;
arrayLengthFromUniform.ubo_binding = {0, kBufferLengthBufferSlot};
+ tint::msl::writer::Bindings bindings;
+
for (BindGroupIndex group : IterateBitSet(layout->GetBindGroupLayoutsMask())) {
- const BindGroupLayoutInternalBase::BindingMap& bindingMap =
- layout->GetBindGroupLayout(group)->GetBindingMap();
- for (const auto [bindingNumber, bindingIndex] : bindingMap) {
- const BindingInfo& bindingInfo =
- layout->GetBindGroupLayout(group)->GetBindingInfo(bindingIndex);
+ const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
- if (!(bindingInfo.visibility & StageBit(stage))) {
- continue;
- }
+ for (const auto& [binding, bindingInfo] : moduleBindingInfo[group]) {
+ tint::BindingPoint srcBindingPoint{static_cast<uint32_t>(group),
+ static_cast<uint32_t>(binding)};
- uint32_t shaderIndex = layout->GetBindingIndexInfo(stage)[group][bindingIndex];
+ BindingIndex bindingIndex = bgl->GetBindingIndex(binding);
+ auto& bindingIndexInfo = layout->GetBindingIndexInfo(stage)[group];
+ uint32_t shaderIndex = bindingIndexInfo[bindingIndex];
- BindingPoint srcBindingPoint{static_cast<uint32_t>(group),
- static_cast<uint32_t>(bindingNumber)};
- BindingPoint dstBindingPoint{0, shaderIndex};
- if (srcBindingPoint != dstBindingPoint) {
- bindingRemapper.binding_points.emplace(srcBindingPoint, dstBindingPoint);
- }
+ tint::BindingPoint dstBindingPoint{0, shaderIndex};
// Use the ShaderIndex as the indices for the buffer size lookups in the array length
// uniform transform. This is used to compute the size of variable length arrays in
@@ -157,6 +149,59 @@
arrayLengthFromUniform.bindpoint_to_size_index.emplace(dstBindingPoint,
dstBindingPoint.binding);
}
+
+ switch (bindingInfo.bindingType) {
+ case BindingInfoType::Buffer:
+ switch (bindingInfo.buffer.type) {
+ case wgpu::BufferBindingType::Uniform:
+ bindings.uniform.emplace(
+ srcBindingPoint,
+ tint::msl::writer::binding::Uniform{dstBindingPoint.binding});
+ break;
+ case kInternalStorageBufferBinding:
+ case wgpu::BufferBindingType::Storage:
+ case wgpu::BufferBindingType::ReadOnlyStorage:
+ bindings.storage.emplace(
+ srcBindingPoint,
+ tint::msl::writer::binding::Storage{dstBindingPoint.binding});
+ break;
+ case wgpu::BufferBindingType::Undefined:
+ DAWN_UNREACHABLE();
+ break;
+ }
+ break;
+ case BindingInfoType::Sampler:
+ bindings.sampler.emplace(srcBindingPoint, tint::msl::writer::binding::Sampler{
+ dstBindingPoint.binding});
+ break;
+ case BindingInfoType::Texture:
+ bindings.texture.emplace(srcBindingPoint, tint::msl::writer::binding::Texture{
+ dstBindingPoint.binding});
+ break;
+ case BindingInfoType::StorageTexture:
+ bindings.storage_texture.emplace(
+ srcBindingPoint,
+ tint::msl::writer::binding::StorageTexture{dstBindingPoint.binding});
+ break;
+ case BindingInfoType::ExternalTexture: {
+ const auto& etBindingMap = bgl->GetExternalTextureBindingExpansionMap();
+ const auto& expansion = etBindingMap.find(binding);
+ DAWN_ASSERT(expansion != etBindingMap.end());
+
+ const auto& bindingExpansion = expansion->second;
+ tint::msl::writer::binding::BindingInfo plane0{
+ static_cast<uint32_t>(shaderIndex)};
+ tint::msl::writer::binding::BindingInfo plane1{
+ bindingIndexInfo[bgl->GetBindingIndex(bindingExpansion.plane1)]};
+ tint::msl::writer::binding::BindingInfo metadata{
+ bindingIndexInfo[bgl->GetBindingIndex(bindingExpansion.params)]};
+
+ bindings.external_texture.emplace(
+ srcBindingPoint,
+ tint::msl::writer::binding::ExternalTexture{metadata, plane0, plane1});
+ break;
+ }
+ }
}
}
@@ -170,11 +215,12 @@
uint32_t metalIndex = renderPipeline->GetMtlVertexBufferIndex(slot);
// Tell Tint to map (kPullingBufferBindingSet, slot) to this MSL buffer index.
- BindingPoint srcBindingPoint{static_cast<uint32_t>(kPullingBufferBindingSet),
- static_cast<uint8_t>(slot)};
- BindingPoint dstBindingPoint{0, metalIndex};
+ tint::BindingPoint srcBindingPoint{static_cast<uint32_t>(kPullingBufferBindingSet),
+ static_cast<uint8_t>(slot)};
+ tint::BindingPoint dstBindingPoint{0, metalIndex};
if (srcBindingPoint != dstBindingPoint) {
- bindingRemapper.binding_points.emplace(srcBindingPoint, dstBindingPoint);
+ bindings.storage.emplace(
+ srcBindingPoint, tint::msl::writer::binding::Storage{dstBindingPoint.binding});
}
// Use the ShaderIndex as the indices for the buffer size lookups in the array
@@ -219,9 +265,8 @@
stage == SingleShaderStage::Vertex &&
renderPipeline->GetPrimitiveTopology() == wgpu::PrimitiveTopology::PointList;
req.tintOptions.array_length_from_uniform = std::move(arrayLengthFromUniform);
- req.tintOptions.binding_remapper_options = std::move(bindingRemapper);
- req.tintOptions.external_texture_options = BuildExternalTextureTransformBindings(layout);
req.tintOptions.pixel_local_options = std::move(pixelLocal);
+ req.tintOptions.bindings = std::move(bindings);
const CombinedLimits& limits = device->GetLimits();
req.limits = LimitsForCompilationRequest::Create(limits.v1);
@@ -352,8 +397,10 @@
}
CacheResult<MslCompilation> mslCompilation;
- DAWN_TRY_ASSIGN(mslCompilation, TranslateToMSL(GetDevice(), programmableStage, stage, layout,
- out, sampleMask, renderPipeline));
+ DAWN_TRY_ASSIGN(mslCompilation,
+ TranslateToMSL(GetDevice(), programmableStage, stage, layout, out, sampleMask,
+ renderPipeline, GetEntryPoint(entryPointName).bindings));
+
out->needsStorageBufferLength = mslCompilation->needsStorageBufferLength;
out->workgroupAllocations = std::move(mslCompilation->workgroupAllocations);
out->localWorkgroupSize = MTLSizeMake(mslCompilation->localWorkgroupSize.width,
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..b570ed4c 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..2b379e5c 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..86ee6be 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_builder.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..db4df49 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_builder.cc",
"options.cc",
"printer_support.cc",
],
hdrs = [
+ "option_builder.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..477eab8 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_builder.cc
+ lang/msl/writer/common/option_builder.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..3c52ebf 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_builder.cc",
+ "option_builder.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/msl/writer/common/option_builder.cc b/src/tint/lang/msl/writer/common/option_builder.cc
new file mode 100644
index 0000000..da45354
--- /dev/null
+++ b/src/tint/lang/msl/writer/common/option_builder.cc
@@ -0,0 +1,222 @@
+/// 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/common/option_builder.h"
+
+#include "src/tint/utils/containers/hashset.h"
+
+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{};
+
+ 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
+ // possible for multiple entry points to use the remapper at the same time. If the pair doesn't
+ // match, then we report an error about a duplicate binding point.
+
+ auto wgsl_seen = [&diagnostics, &seen_wgsl_bindings](const tint::BindingPoint& src,
+ const binding::BindingInfo& dst) -> bool {
+ if (auto binding = seen_wgsl_bindings.Find(src)) {
+ if (*binding != dst) {
+ std::stringstream str;
+ str << "found duplicate WGSL binding point: " << src;
+
+ diagnostics.add_error(diag::System::Writer, str.str());
+ return true;
+ }
+ }
+ seen_wgsl_bindings.Add(src, dst);
+ return false;
+ };
+
+ 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 MSL binding point: [binding: " << src.binding << "]";
+ diagnostics.add_error(diag::System::Writer, str.str());
+ return true;
+ }
+ }
+ map.Add(src, dst);
+ return false;
+ };
+
+ 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;
+
+ if (wgsl_seen(src_binding, dst_binding)) {
+ return false;
+ }
+
+ if (msl_seen(map, dst_binding, src_binding)) {
+ return false;
+ }
+ }
+ return true;
+ };
+
+ // 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(seen_msl_buffer_bindings, options.bindings.storage)) {
+ diagnostics.add_note(diag::System::Writer, "when processing storage", {});
+ return false;
+ }
+
+ // 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(seen_msl_texture_bindings, options.bindings.storage_texture)) {
+ diagnostics.add_note(diag::System::Writer, "when processing storage_texture", {});
+ return false;
+ }
+
+ for (const auto& it : options.bindings.external_texture) {
+ const auto& src_binding = it.first;
+ const auto& plane0 = it.second.plane0;
+ const auto& plane1 = it.second.plane1;
+ const auto& metadata = it.second.metadata;
+
+ // Validate with the actual source regardless of what the remapper will do
+ if (wgsl_seen(src_binding, plane0)) {
+ diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
+ return false;
+ }
+
+ // 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 (msl_seen(seen_msl_texture_bindings, plane1, src_binding)) {
+ diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
+ return false;
+ }
+ // 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;
+ }
+ }
+
+ return true;
+}
+
+// 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 MSL
+// (binding) in the `uniform`, `storage`, `texture`, and `sampler` arrays.
+void PopulateRemapperAndMultiplanarOptions(const Options& options,
+ RemapperData& remapper_data,
+ ExternalTextureOptions& external_texture) {
+ auto create_remappings = [&remapper_data](const auto& hsh) {
+ for (const auto& it : hsh) {
+ const BindingPoint& src_binding_point = it.first;
+ const binding::Uniform& dst_binding_point = it.second;
+
+ // 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{0, dst_binding_point.binding});
+ }
+ };
+
+ create_remappings(options.bindings.uniform);
+ create_remappings(options.bindings.storage);
+ create_remappings(options.bindings.texture);
+ create_remappings(options.bindings.storage_texture);
+ create_remappings(options.bindings.sampler);
+
+ // 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{0, plane0.binding};
+ BindingPoint plane1_binding_point{0, plane1.binding};
+ BindingPoint metadata_binding_point{0, metadata.binding};
+
+ // Use the re-bound msl plane0 value for the lookup key.
+ external_texture.bindings_map.emplace(
+ 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 MSL do not need to be re-bound.
+ if (src_binding_point == plane0_binding_point) {
+ continue;
+ }
+
+ remapper_data.emplace(src_binding_point, plane0_binding_point);
+ }
+}
+
+} // namespace tint::msl::writer
+
+namespace std {
+
+/// 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::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::msl::writer::binding::BindingInfo& info) const {
+ return tint::Hash(info.binding);
+ }
+};
+
+} // namespace std
diff --git a/src/tint/lang/msl/writer/common/option_builder.h b/src/tint/lang/msl/writer/common/option_builder.h
new file mode 100644
index 0000000..4f63ad4
--- /dev/null
+++ b/src/tint/lang/msl/writer/common/option_builder.h
@@ -0,0 +1,59 @@
+/// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_BUILDER_H_
+#define SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_BUILDER_H_
+
+#include <unordered_map>
+
+#include "src/tint/api/common/binding_point.h"
+#include "src/tint/api/options/external_texture.h"
+#include "src/tint/lang/msl/writer/common/options.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+
+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);
+
+/// Populates data from the writer options for the remapper and external texture.
+/// @param options the writer options
+/// @param remapper_data where to put the remapper data
+/// @param external_texture where to store the external texture options
+/// Note, these are populated together because there are dependencies between the two types of data.
+void PopulateRemapperAndMultiplanarOptions(const Options& options,
+ RemapperData& remapper_data,
+ ExternalTextureOptions& external_texture);
+
+} // namespace tint::msl::writer
+
+#endif // SRC_TINT_LANG_MSL_WRITER_COMMON_OPTION_BUILDER_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..e720b99 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_builder.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/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/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
