[dawn][vk] Set up the metadata buffer for Tint's resource_binding.
Tint's support for `resource_binding` is done with a transform that
needs to know where the metadata buffer will be. Add that information in
the options set up in vulkan::ShaderModule.
Add tests that the metadata buffer is correctly set up by checking that
WGSL's arrayLength on `resource_binding` works correctly.
Bug: 435317394, 442483669
Change-Id: Ie90d857047f621bac7fc0c58c21b39e4683fc7a3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/259874
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
diff --git a/src/dawn/native/BindGroup.cpp b/src/dawn/native/BindGroup.cpp
index 3626965..860dbbd 100644
--- a/src/dawn/native/BindGroup.cpp
+++ b/src/dawn/native/BindGroup.cpp
@@ -984,6 +984,12 @@
return {};
}
+DynamicArrayState* BindGroupBase::GetDynamicArray() const {
+ DAWN_ASSERT(!IsError());
+ DAWN_ASSERT(HasDynamicArray());
+ return mDynamicArray.get();
+}
+
MaybeError BindGroupBase::ValidateDestroy() const {
DAWN_TRY(GetDevice()->ValidateObject(this));
diff --git a/src/dawn/native/BindGroup.h b/src/dawn/native/BindGroup.h
index 77bcafc..9144ffc 100644
--- a/src/dawn/native/BindGroup.h
+++ b/src/dawn/native/BindGroup.h
@@ -91,6 +91,7 @@
bool HasDynamicArray() const;
ityp::span<BindingIndex, const Ref<TextureViewBase>> GetDynamicArrayBindings() const;
MaybeError ValidateCanUseOnQueueNow() const;
+ DynamicArrayState* GetDynamicArray() const;
protected:
// To save memory, the size of a bind group is dynamically determined and the bind group is
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index 61cb081..8a03626 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -222,6 +222,40 @@
}
}
+ // Add options for dynamic binding arrays. They need remapping like all regular bindings but
+ // also need to give information about additional bindings for the metadata buffer and the
+ // default bindings.
+ tint::ResourceBindingConfig resourceBindingConfig;
+ for (BindGroupIndex group : layout->GetBindGroupLayoutsMask()) {
+ const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
+ if (!bgl->HasDynamicArray()) {
+ continue;
+ }
+
+ tint::BindingPoint wgslDynamicArrayBindPoint = {
+ .group = uint32_t(group), .binding = uint32_t(bgl->GetAPIDynamicArrayStart())};
+ tint::BindingPoint remappedDynamicArrayBindPoint = {
+ .group = uint32_t(group),
+ .binding = uint32_t(bgl->GetDynamicArrayStart()),
+ };
+ tint::BindingPoint metadataBindPoint = {
+ .group = uint32_t(group),
+ .binding = uint32_t(bgl->GetDynamicArrayMetadataBinding()),
+ };
+
+ // TODO(https://crbug.com/442483669): This uses the texture binding remapper support to
+ // remap a `resource_binding`. It is a hack until Tint adds support for `resource_binding`
+ // to the binding remapper.
+ bindings.texture.emplace(wgslDynamicArrayBindPoint, remappedDynamicArrayBindPoint);
+
+ // The resourceBindingConfig only uses remapped bind points.
+ resourceBindingConfig.bindings[remappedDynamicArrayBindPoint] = {
+ .storage_buffer_binding = metadataBindPoint,
+ // TODO(https://crbug.com/435317394): Support for all the resource types.
+ .default_binding_type_order = {},
+ };
+ }
+
const bool hasInputAttachment = !bindings.input_attachment.empty();
SpirvCompilationRequest req = {};
@@ -253,6 +287,7 @@
req.tintOptions.use_storage_input_output_16 =
GetDevice()->IsToggleEnabled(Toggle::VulkanUseStorageInputOutput16);
req.tintOptions.bindings = std::move(bindings);
+ req.tintOptions.resource_binding = std::move(resourceBindingConfig);
req.tintOptions.disable_image_robustness =
GetDevice()->IsToggleEnabled(Toggle::VulkanUseImageRobustAccess2);
// Currently we can disable index clamping on all runtime-sized arrays in Tint robustness
diff --git a/src/dawn/tests/end2end/BindingArrayTests.cpp b/src/dawn/tests/end2end/BindingArrayTests.cpp
index 2ac4811..156cb00 100644
--- a/src/dawn/tests/end2end/BindingArrayTests.cpp
+++ b/src/dawn/tests/end2end/BindingArrayTests.cpp
@@ -810,6 +810,72 @@
// - Check that a newly created resource that's pinned samples as zeroes.
// - Likewise for a texture written to, then discarded with a render pass.
+// Test that the WGSL `arrayLength` builtin on dynamic binding arrays returns the correct length.
+TEST_P(DynamicBindingArrayTests, ArrayLengthBuiltin) {
+ // Create a compute pipeline that returns the array length of the dynamic binding arrays.
+ // One of them has a static binding as well so as to check that it doesn't mess up the
+ // computation of the array length.
+ wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+
+ @group(0) @binding(0) var<storage, read_write> result : array<u32, 2>;
+ @group(0) @binding(1) var firstBindings : resource_binding;
+ @group(1) @binding(0) var secondBindings : resource_binding;
+
+ @compute @workgroup_size(1) fn getArrayLengths() {
+ // Force the defaulted layout to wgpu::DynamicBindingKind::SampledTexture
+ _ = hasBinding<texture_2d<f32>>(firstBindings, 0);
+ _ = hasBinding<texture_2d<f32>>(secondBindings, 0);
+
+ result[0] = arrayLength(firstBindings);
+ result[1] = arrayLength(secondBindings);
+ }
+ )");
+ wgpu::ComputePipelineDescriptor csDesc = {.compute = {
+ .module = module,
+ }};
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+ // Create the dynamic binding arrays and fetch their array length in a buffer.
+ wgpu::BufferDescriptor bDesc = {
+ .usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc,
+ .size = 2 * sizeof(uint32_t),
+ };
+ wgpu::Buffer resultBuffer = device.CreateBuffer(&bDesc);
+
+ wgpu::BindGroup bg0 = MakeBindGroup(pipeline.GetBindGroupLayout(0), 17, {{0, resultBuffer}});
+ wgpu::BindGroup bg1 = MakeBindGroup(pipeline.GetBindGroupLayout(1), 3, {});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetBindGroup(0, bg0);
+ pass.SetBindGroup(1, bg1);
+ pass.SetPipeline(pipeline);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+
+ wgpu::CommandBuffer commands = encoder.Finish();
+ device.GetQueue().Submit(1, &commands);
+
+ // The result buffer should contain the lengths of the dynamic binding arrays.
+ EXPECT_BUFFER_U32_EQ(17, resultBuffer, 0);
+ EXPECT_BUFFER_U32_EQ(3, resultBuffer, 4);
+}
+
+// TODO(https://crbug.com/435317394): Add tests that texture pinning / unpinning is reflected in the
+// availability of the binding in the shader. This can be done with a compute shader that loops over
+// [0, arrayLength(&resource_binding)) and returns hasBinding<T>(i) in a result storage buffer.
+// - Test adding one texture binding in the middle of a dynamic binding array, unpinned the pinned
+// - Test adding one texture binding multiple times, pinning it then unpinning it.
+// - Test adding multiple textures and the same texture twice (pinned). then unpinning one of them
+// - Test adding a texture, pinning it, then destroying it.
+// - Test adding the same texture to multiple dynamic binding arrays, pinning it, then destroying
+// one of the arrays.
+// - Add TODO to test with binding array updates in the future.
+// - Test for each possible sampled type, a shader that check hasBinding on every possible sampled
+// type as well (to know that we correctly pass in the type ID).
+// - Start at other dynamicArrayStart
+
DAWN_INSTANTIATE_TEST(DynamicBindingArrayTests, D3D12Backend(), MetalBackend(), VulkanBackend());
} // anonymous namespace
