[dawn][native] Update typeId metadata for dynamic binding arrays.
This metadata is used inside shaders to know what to return for
hasBinding<T>(resource_binding, iu32) for the various kinds of T.
This CL adds all the mechanisms to track and update the typeIds but only
makes the difference between "empty" and "texture_2d<f32>".
- In Textures, add a map of all the (DynamicArrayState, slot) they are
referenced in, and maintain it as DynamicArrayStates are created and
destroyed. (this map contains WeakRefs to the DynamicArrayState to
break what would otherwise be a ref cycle),
- Add support for Hash to WeakRef for the above map.
- Notify the DynamicArrayState on texture.Pin/Unpin so that it can
track the "pinned" status of each slot.
- In DynamicArrayState maintain the tint::ResourceType for each slot
and return a list of updates to the metadata that backends must
perform before using the dynamic binding array.
- In the Vulkan backend, implement this update of the binding array
metadata.
- Fix an issue in Vulkan where setting a destroyed texture in the
dynamic binding array at creation, would cause a VK_NULL_HANDLE to be
written in vkWriteDescriptorSets, which is invalid.
- Add a bunch of tests that hasBinding works correctly to exercise the
newly added code paths.
Bug: 435317394
Change-Id: Ie5a4568f033cb3f8d0d7dac29a9575514564a32e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/259875
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Auto-Submit: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
diff --git a/src/dawn/common/WeakRef.h b/src/dawn/common/WeakRef.h
index 5322b59..16adf53 100644
--- a/src/dawn/common/WeakRef.h
+++ b/src/dawn/common/WeakRef.h
@@ -118,6 +118,8 @@
// Friend is needed so that hashing can access the mData pointer which is always valid.
template <typename U, typename H>
friend H AbslHashValue(H h, const WeakRef<U>& v);
+ template <typename U>
+ friend size_t Hash(const WeakRef<U>& v);
// Constructor from data should only be allowed via the GetWeakRef function.
explicit WeakRef(detail::WeakRefSupportBase* data) : mData(data->mData) {}
@@ -130,6 +132,11 @@
return H::combine(std::move(h), v.mData.Get());
}
+template <typename T>
+size_t Hash(const WeakRef<T>& v) {
+ return Hash(v.mData.Get());
+}
+
} // namespace dawn
#endif // SRC_DAWN_COMMON_WEAKREFCOUNTED_H_
diff --git a/src/dawn/native/BindGroup.cpp b/src/dawn/native/BindGroup.cpp
index 860dbbd..a61efcc 100644
--- a/src/dawn/native/BindGroup.cpp
+++ b/src/dawn/native/BindGroup.cpp
@@ -815,8 +815,7 @@
// Gather dynamic binding entries in a second loop to put the handling off the critical path.
if (auto* dynamic = descriptor.Get<BindGroupDynamicBindingArray>()) {
- mDynamicArray =
- std::make_unique<DynamicArrayState>(BindingIndex(dynamic->dynamicArraySize));
+ mDynamicArray = AcquireRef(new DynamicArrayState(BindingIndex(dynamic->dynamicArraySize)));
DAWN_TRY(mDynamicArray->Initialize(GetDevice()));
for (uint32_t i = 0; i < descriptor->entryCount; ++i) {
@@ -987,7 +986,7 @@
DynamicArrayState* BindGroupBase::GetDynamicArray() const {
DAWN_ASSERT(!IsError());
DAWN_ASSERT(HasDynamicArray());
- return mDynamicArray.get();
+ return mDynamicArray.Get();
}
MaybeError BindGroupBase::ValidateDestroy() const {
diff --git a/src/dawn/native/BindGroup.h b/src/dawn/native/BindGroup.h
index 9144ffc..e98e948 100644
--- a/src/dawn/native/BindGroup.h
+++ b/src/dawn/native/BindGroup.h
@@ -29,7 +29,6 @@
#define SRC_DAWN_NATIVE_BINDGROUP_H_
#include <array>
-#include <memory>
#include <vector>
#include "dawn/common/Constants.h"
@@ -133,8 +132,9 @@
std::vector<Ref<ExternalTextureBase>> mBoundExternalTextures;
// The dynamic array is separate so as to not bloat the size and destructor of bind groups
- // without them.
- std::unique_ptr<DynamicArrayState> mDynamicArray;
+ // without them. Note that this is the only persistent owning Ref. DynamicArray is a RefCounted
+ // only so WeakRef to it can be created.
+ Ref<DynamicArrayState> mDynamicArray;
};
} // namespace dawn::native
diff --git a/src/dawn/native/DynamicArrayState.cpp b/src/dawn/native/DynamicArrayState.cpp
index d5902d6..64bc9b7 100644
--- a/src/dawn/native/DynamicArrayState.cpp
+++ b/src/dawn/native/DynamicArrayState.cpp
@@ -27,34 +27,82 @@
#include "dawn/native/DynamicArrayState.h"
+#include "dawn/common/Enumerator.h"
#include "dawn/native/Buffer.h"
#include "dawn/native/Device.h"
+#include "tint/api/common/resource_type.h"
namespace dawn::native {
+namespace {
+
+tint::ResourceType ComputeTypeId(const TextureViewBase* view) {
+ // TODO(https://crbug.com/439522242): Add handling of all sampled texture types.
+ if (view == nullptr) {
+ return tint::ResourceType::kEmpty;
+ }
+
+ // TODO(https://crbug.com/435317394): Add support for all the other types that can be in
+ // DynamicBindingKind::SampledTexture. Hardcode to texture_2d<f32> for now.
+ return tint::ResourceType::kTexture2d_f32;
+}
+
+} // anonymous namespace
+
DynamicArrayState::DynamicArrayState(BindingIndex size) {
mBindings.resize(size);
+
+ DAWN_ASSERT(ComputeTypeId(nullptr) == BindingState{}.typeId);
+ mBindingState.resize(size);
}
MaybeError DynamicArrayState::Initialize(DeviceBase* device) {
+ // TODO(https://crbug.com/435317394): Default bindings will be included in mBindings in the
+ // future such that we should use the dynamicArraySize passed in the BindGroup creation instead
+ // of the size of mBindings.
+ uint32_t metadataArrayLength = uint32_t(mBindings.size());
+
// Create a storage buffer that will hold the shader-visible metadata for the dynamic array.
BufferDescriptor metadataDesc{
.label = "binding array metadata",
.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopyDst,
- .size = 4,
+ .size = sizeof(uint32_t) * (metadataArrayLength + 1),
.mappedAtCreation = true,
};
DAWN_TRY_ASSIGN(mMetadataBuffer, device->CreateBuffer(&metadataDesc));
- // TODO(https://crbug.com/439522242): For now it only contains the size but we also need to add
- // type information for each entry in the future.
+ // Initialize the metadata buffer with the arrayLength and a bunch of zeroes that correspond to
+ // empty entries.
+ DAWN_ASSERT(uint32_t(tint::ResourceType::kEmpty) == 0);
+ // TODO(https://crbug.com/435317394): We could rely on zero initialization if it is enabled, and
+ // also apply the initial dirty bindings in this mapping instead of one the first use of the
+ // dynamic binding array.
uint32_t* data = static_cast<uint32_t*>(mMetadataBuffer->GetMappedRange(0, metadataDesc.size));
- *data = uint32_t(mBindings.size());
+ *data = metadataArrayLength;
+ memset(data + 1, 0, metadataDesc.size - sizeof(uint32_t));
DAWN_TRY(mMetadataBuffer->Unmap());
return {};
}
+void DynamicArrayState::Destroy() {
+ DAWN_ASSERT(!mDestroyed);
+
+ for (auto [i, view] : Enumerate(mBindings)) {
+ if (view != nullptr) {
+ view->GetTexture()->RemoveDynamicArraySlot(this, i);
+ }
+ }
+
+ mBindings.clear();
+ mBindingState.clear();
+ mDirtyBindings.clear();
+ mMetadataBuffer->Destroy();
+ mMetadataBuffer = nullptr;
+
+ mDestroyed = true;
+}
+
BindingIndex DynamicArrayState::GetSize() const {
DAWN_ASSERT(!mDestroyed);
return mBindings.size();
@@ -76,15 +124,81 @@
void DynamicArrayState::Update(BindingIndex i, TextureViewBase* view) {
DAWN_ASSERT(!mDestroyed);
+ if (mBindings[i] == view) {
+ return;
+ }
+
+ // Update the mBindings slot but also the mapping to the slot that are stored in the textures.
+ if (mBindings[i] != nullptr) {
+ mBindings[i]->GetTexture()->RemoveDynamicArraySlot(this, i);
+ }
+ if (view != nullptr) {
+ view->GetTexture()->AddDynamicArraySlot(this, i);
+ }
mBindings[i] = view;
+
+ // Update the mBindingState with information for the updated binding.
+ tint::ResourceType typeId = ComputeTypeId(view);
+ bool pinned = view != nullptr && view->GetTexture()->HasPinnedUsage();
+
+ BindingState& state = mBindingState[i];
+ if (state.typeId != typeId || state.pinned != pinned) {
+ state.typeId = typeId;
+ state.pinned = pinned;
+ MarkStateDirty(i);
+ }
}
-void DynamicArrayState::Destroy() {
+void DynamicArrayState::OnPinned(BindingIndex i, TextureBase* texture) {
DAWN_ASSERT(!mDestroyed);
- mBindings.clear();
- mMetadataBuffer->Destroy();
- mMetadataBuffer = nullptr;
- mDestroyed = true;
+ DAWN_ASSERT(mBindings[i] != nullptr);
+ DAWN_ASSERT(mBindings[i]->GetTexture() == texture);
+ DAWN_ASSERT(!mBindingState[i].pinned);
+ mBindingState[i].pinned = true;
+ MarkStateDirty(i);
+}
+
+void DynamicArrayState::OnUnpinned(BindingIndex i, TextureBase* texture) {
+ DAWN_ASSERT(!mDestroyed);
+ DAWN_ASSERT(mBindings[i] != nullptr);
+ DAWN_ASSERT(mBindings[i]->GetTexture() == texture);
+ DAWN_ASSERT(mBindingState[i].pinned);
+ mBindingState[i].pinned = false;
+ MarkStateDirty(i);
+}
+
+bool DynamicArrayState::HasDirtyBindings() const {
+ return !mDirtyBindings.empty();
+}
+
+std::vector<DynamicArrayState::BindingStateUpdate> DynamicArrayState::AcquireDirtyBindingUpdates() {
+ DAWN_ASSERT(!mDestroyed);
+
+ std::vector<BindingStateUpdate> updates;
+ for (BindingIndex dirtyIndex : mDirtyBindings) {
+ DAWN_ASSERT(mBindingState[dirtyIndex].dirty);
+ mBindingState[dirtyIndex].dirty = false;
+
+ tint::ResourceType effectiveType = mBindingState[dirtyIndex].pinned
+ ? mBindingState[dirtyIndex].typeId
+ : tint::ResourceType::kEmpty;
+
+ size_t offset = sizeof(uint32_t) * (uint32_t(dirtyIndex) + 1);
+ updates.push_back({
+ .offset = uint32_t(offset),
+ .data = uint32_t(effectiveType),
+ });
+ }
+ mDirtyBindings.clear();
+
+ return updates;
+}
+
+void DynamicArrayState::MarkStateDirty(BindingIndex i) {
+ if (!mBindingState[i].dirty) {
+ mDirtyBindings.push_back(i);
+ mBindingState[i].dirty = true;
+ }
}
} // namespace dawn::native
diff --git a/src/dawn/native/DynamicArrayState.h b/src/dawn/native/DynamicArrayState.h
index 1ddeae7..b330949 100644
--- a/src/dawn/native/DynamicArrayState.h
+++ b/src/dawn/native/DynamicArrayState.h
@@ -31,6 +31,8 @@
#include <vector>
#include "dawn/common/Ref.h"
+#include "dawn/common/RefCounted.h"
+#include "dawn/common/WeakRefSupport.h"
#include "dawn/common/ityp_span.h"
#include "dawn/common/ityp_vector.h"
#include "dawn/native/Error.h"
@@ -39,30 +41,72 @@
#include "dawn/native/dawn_platform.h"
#include "partition_alloc/pointers/raw_ptr.h"
+namespace tint {
+enum class ResourceType : uint32_t;
+} // namespace tint
+
namespace dawn::native {
// An optional component of a BindGroup that's used to track the resources that are in the dynamic
// binding array part. It helps maintain the metadata buffer that's used in shaders to know if it is
// valid to access an entry of the dynamic binding array with a given type (note that the writing of
// the updates to the buffer are done by the backends).
-class DynamicArrayState {
+//
+// DynamicArrayState has a single strong reference owned by the BindGroup that created it, however
+// all resources contained in the dynamic array need WeakRefs to update it on Pin/Unpin. (They use
+// WeakRef to avoid a reference cycle between the dynamic array and its bindings).
+class DynamicArrayState : public RefCounted, public WeakRefSupport<DynamicArrayState> {
public:
explicit DynamicArrayState(BindingIndex size);
MaybeError Initialize(DeviceBase* device);
+ void Destroy();
BindingIndex GetSize() const;
ityp::span<BindingIndex, const Ref<TextureViewBase>> GetBindings() const;
BufferBase* GetMetadataBuffer() const;
bool IsDestroyed() const;
+ // Methods that mutate the state of bindings in the dynamic array. They keep track of the
+ // necessary metadata buffer updates required for dynamic type checks in the shader to match
+ // what's in the binding array.
void Update(BindingIndex i, TextureViewBase* view);
- void Destroy();
+ void OnPinned(BindingIndex i, TextureBase* texture);
+ void OnUnpinned(BindingIndex i, TextureBase* texture);
+
+ // Returns the various type ids that need to be updated in the metadata buffer before the next
+ // use of the binding array.
+ struct BindingStateUpdate {
+ uint32_t offset;
+ uint32_t data;
+ };
+ bool HasDirtyBindings() const;
+ std::vector<BindingStateUpdate> AcquireDirtyBindingUpdates();
private:
bool mDestroyed = false;
+
ityp::vector<BindingIndex, Ref<TextureViewBase>> mBindings;
+ // Buffer that contains a WGSL metadata struct of the following shape:
+ //
+ // struct Metadata {
+ // arrayLength: u32, // Doesn't include the default bindings
+ // bindings: array<u32>, // `arrayLength` entries
+ // }
Ref<BufferBase> mMetadataBuffer;
+
+ struct BindingState {
+ // Matches the value of the Tint enum for type IDs but kept as u32 to keep usage of Tint
+ // headers local.
+ tint::ResourceType typeId = tint::ResourceType(0);
+ bool dirty = false;
+ bool pinned = false;
+ };
+ ityp::vector<BindingIndex, BindingState> mBindingState;
+
+ // The list of bindings that need to be updated before the next use of the dynamic array.
+ std::vector<BindingIndex> mDirtyBindings;
+ void MarkStateDirty(BindingIndex i);
};
} // namespace dawn::native
diff --git a/src/dawn/native/Texture.cpp b/src/dawn/native/Texture.cpp
index d4ce94b..e61fc0f 100644
--- a/src/dawn/native/Texture.cpp
+++ b/src/dawn/native/Texture.cpp
@@ -41,6 +41,7 @@
#include "dawn/native/ChainUtils.h"
#include "dawn/native/CommandValidation.h"
#include "dawn/native/Device.h"
+#include "dawn/native/DynamicArrayState.h"
#include "dawn/native/EnumMaskIterator.h"
#include "dawn/native/ObjectType_autogen.h"
#include "dawn/native/PassResourceUsage.h"
@@ -1587,6 +1588,22 @@
// Update the frontend state.
mPinnedUsage = usage;
+
+ // Call OnPinned for each of the slots. We would like to prune the entries to now destroyed
+ // DynamicArrayStates using the `it = set.erase(it)` std:: idiom, but that's not possible with
+ // absl::flat_hash_set. Instead track a list of entries to prune and do it in a second pass.
+ std::vector<DynamicArraySlot> slotsToPrune;
+ for (const auto& slot : mDynamicArraySlots) {
+ if (Ref<DynamicArrayState> dynamicArray = slot.dynamicArray.Promote()) {
+ dynamicArray->OnPinned(slot.slot, this);
+ } else {
+ slotsToPrune.push_back(slot);
+ }
+ }
+ for (const auto& slot : slotsToPrune) {
+ mDynamicArraySlots.erase(slot);
+ }
+
return {};
}
@@ -1613,6 +1630,48 @@
// Update the frontend state.
mPinnedUsage = wgpu::TextureUsage::None;
+
+ // Call OnUnpinned for each of the slots. We would like to prune the entries to now destroyed
+ // DynamicArrayStates using the `it = set.erase(it)` std:: idiom, but that's not possible with
+ // absl::flat_hash_set. Instead track a list of entries to prune and do it in a second pass.
+ std::vector<DynamicArraySlot> slotsToPrune;
+ for (const auto& slot : mDynamicArraySlots) {
+ if (Ref<DynamicArrayState> dynamicArray = slot.dynamicArray.Promote()) {
+ dynamicArray->OnUnpinned(slot.slot, this);
+ } else {
+ slotsToPrune.push_back(slot);
+ }
+ }
+ for (const auto& slot : slotsToPrune) {
+ mDynamicArraySlots.erase(slot);
+ }
+}
+
+void TextureBase::AddDynamicArraySlot(DynamicArrayState* dynamicArray, BindingIndex i) {
+ DAWN_ASSERT(!IsError());
+ // Note that this can be called after the texture is destroyed.
+ DynamicArraySlot slot = {dynamicArray, i};
+ auto [_, inserted] = mDynamicArraySlots.insert(slot);
+ DAWN_ASSERT(inserted);
+}
+
+void TextureBase::RemoveDynamicArraySlot(DynamicArrayState* dynamicArray, BindingIndex i) {
+ DAWN_ASSERT(!IsError());
+ // Note that this can be called after the texture is destroyed.
+ DynamicArraySlot slot = {dynamicArray, i};
+ bool removed = mDynamicArraySlots.erase(slot);
+ DAWN_ASSERT(removed);
+}
+
+size_t TextureBase::DynamicArraySlot::HashFuncs::operator()(const DynamicArraySlot& query) const {
+ size_t hash = 0;
+ HashCombine(&hash, query.dynamicArray, query.slot);
+ return hash;
+}
+
+bool TextureBase::DynamicArraySlot::HashFuncs::operator()(const DynamicArraySlot& a,
+ const DynamicArraySlot& b) const {
+ return std::tie(a.dynamicArray, a.slot) == std::tie(b.dynamicArray, b.slot);
}
void TextureBase::UnpinImpl() {
diff --git a/src/dawn/native/Texture.h b/src/dawn/native/Texture.h
index 4510c9f..d1f8878 100644
--- a/src/dawn/native/Texture.h
+++ b/src/dawn/native/Texture.h
@@ -32,6 +32,7 @@
#include <string>
#include <vector>
+#include "absl/container/flat_hash_set.h"
#include "dawn/common/LRUCache.h"
#include "dawn/common/RefCountedWithExternalCount.h"
#include "dawn/common/WeakRef.h"
@@ -40,15 +41,16 @@
#include "dawn/native/Error.h"
#include "dawn/native/Format.h"
#include "dawn/native/Forward.h"
+#include "dawn/native/IntegerTypes.h"
#include "dawn/native/ObjectBase.h"
#include "dawn/native/SharedTextureMemory.h"
#include "dawn/native/Subresource.h"
-#include "partition_alloc/pointers/raw_ref.h"
-
#include "dawn/native/dawn_platform.h"
+#include "partition_alloc/pointers/raw_ref.h"
namespace dawn::native {
+class DynamicArrayState;
class MemoryDump;
enum class AllowMultiPlanarTextureFormat {
@@ -190,6 +192,8 @@
MaybeError Pin(wgpu::TextureUsage usage);
void Unpin();
+ void AddDynamicArraySlot(DynamicArrayState* dynamicArray, BindingIndex i);
+ void RemoveDynamicArraySlot(DynamicArrayState* dynamicArray, BindingIndex i);
ResultOrError<Ref<TextureViewBase>> CreateView(
const TextureViewDescriptor* descriptor = nullptr);
@@ -281,6 +285,19 @@
LRUCache<TextureViewQuery, Ref<TextureViewBase>, TextureViewCacheFuncs>;
std::unique_ptr<TextureViewCache> mTextureViewCache;
+ // Keep a hash set of the places this texture is bound to in DynamicArrayStates.
+ struct DynamicArraySlot {
+ WeakRef<DynamicArrayState> dynamicArray;
+ BindingIndex slot;
+
+ struct HashFuncs {
+ size_t operator()(const DynamicArraySlot& query) const;
+ bool operator()(const DynamicArraySlot& a, const DynamicArraySlot& b) const;
+ };
+ };
+ absl::flat_hash_set<DynamicArraySlot, DynamicArraySlot::HashFuncs, DynamicArraySlot::HashFuncs>
+ mDynamicArraySlots;
+
// TODO(crbug.com/dawn/845): Use a more optimized data structure to save space
std::vector<bool> mIsSubresourceContentInitializedAtIndex;
};
diff --git a/src/dawn/native/vulkan/BindGroupVk.cpp b/src/dawn/native/vulkan/BindGroupVk.cpp
index c375e289..da70d2d 100644
--- a/src/dawn/native/vulkan/BindGroupVk.cpp
+++ b/src/dawn/native/vulkan/BindGroupVk.cpp
@@ -223,14 +223,22 @@
auto bindings = GetDynamicArrayBindings();
for (auto [i, view] : Enumerate(bindings)) {
- if (view != nullptr) {
- VkDescriptorImageInfo imageWrite = {};
- imageWrite.imageView = ToBackend(view)->GetHandle();
- imageWrite.imageLayout =
- VulkanImageLayout(view->GetFormat(), wgpu::TextureUsage::TextureBinding);
- imageWrites.push_back(imageWrite);
- arrayElements.push_back(uint32_t(i));
+ if (view == nullptr) {
+ continue;
}
+
+ VkImageView handle = ToBackend(view)->GetHandle();
+ if (handle == nullptr) {
+ continue;
+ }
+
+ VkDescriptorImageInfo imageWrite = {
+ .sampler = VkSampler{},
+ .imageView = handle,
+ .imageLayout = VulkanImageLayout(view->GetFormat(), wgpu::TextureUsage::TextureBinding),
+ };
+ imageWrites.push_back(imageWrite);
+ arrayElements.push_back(uint32_t(i));
}
std::vector<VkWriteDescriptorSet> writes;
diff --git a/src/dawn/native/vulkan/CommandBufferVk.cpp b/src/dawn/native/vulkan/CommandBufferVk.cpp
index eba7850..5ff6edd 100644
--- a/src/dawn/native/vulkan/CommandBufferVk.cpp
+++ b/src/dawn/native/vulkan/CommandBufferVk.cpp
@@ -31,10 +31,12 @@
#include <limits>
#include <vector>
+#include "dawn/common/Enumerator.h"
#include "dawn/native/BindGroupTracker.h"
#include "dawn/native/CommandEncoder.h"
#include "dawn/native/CommandValidation.h"
#include "dawn/native/Commands.h"
+#include "dawn/native/DynamicArrayState.h"
#include "dawn/native/DynamicUploader.h"
#include "dawn/native/EnumMaskIterator.h"
#include "dawn/native/ImmediateConstantsTracker.h"
@@ -226,11 +228,67 @@
}
};
+// Updates a dynamic array metadata buffer by scheduling a copy for each u32 that needs to be
+// updated.
+// TODO(https://crbug.com/435317394): If we had a way to Dawn reentrantly now, we could use a
+// compute shader to dispatch the updates instead of individual copies for each update, and move
+// that logic in the frontend to share it between backends. (also a single dispatch could update
+// multiple metadata buffers potentially).
+MaybeError UpdateDynamicArrayBindings(Device* device,
+ CommandRecordingContext* recordingContext,
+ DynamicArrayState* dynamicArray) {
+ std::vector<DynamicArrayState::BindingStateUpdate> updates =
+ dynamicArray->AcquireDirtyBindingUpdates();
+
+ // Allocate enough space for all the data to modify and schedule the copies.
+ DAWN_TRY(device->GetDynamicUploader()->WithUploadReservation(
+ sizeof(uint32_t) * updates.size(), kCopyBufferToBufferOffsetAlignment,
+ [&](UploadReservation reservation) -> MaybeError {
+ uint32_t* stagedData = static_cast<uint32_t*>(reservation.mappedPointer);
+
+ // The metadata buffer will be copied to.
+ Buffer* metadataBuffer = ToBackend(dynamicArray->GetMetadataBuffer());
+ DAWN_ASSERT(metadataBuffer->IsInitialized());
+ metadataBuffer->TransitionUsageNow(recordingContext, wgpu::BufferUsage::CopyDst);
+
+ // Prepare the copies.
+ std::vector<VkBufferCopy> copies(updates.size());
+ for (auto [i, update] : Enumerate(updates)) {
+ stagedData[i] = update.data;
+
+ VkBufferCopy copy{
+ .srcOffset = reservation.offsetInBuffer + i * sizeof(uint32_t),
+ .dstOffset = update.offset,
+ .size = sizeof(uint32_t),
+ };
+ copies[i] = copy;
+ }
+
+ // Enqueue the copy commands all at once.
+ device->fn.CmdCopyBuffer(recordingContext->commandBuffer,
+ ToBackend(reservation.buffer)->GetHandle(),
+ metadataBuffer->GetHandle(), copies.size(), copies.data());
+ return {};
+ }));
+
+ return {};
+}
+
// Records the necessary barriers for a synchronization scope using the resource usage
// data pre-computed in the frontend. Also performs lazy initialization if required.
-MaybeError TransitionAndClearForSyncScope(Device* device,
- CommandRecordingContext* recordingContext,
- const SyncScopeResourceUsage& scope) {
+MaybeError PrepareResourcesForSyncScope(Device* device,
+ CommandRecordingContext* recordingContext,
+ const SyncScopeResourceUsage& scope) {
+ // Update the dynamic binding array metadata buffers before transitioning resources. The
+ // metadata buffers are part of the resources and will be transitioned to Storage if needed
+ // then.
+ for (BindGroupBase* dynamicArrayBG : scope.dynamicBindingArrays) {
+ if (dynamicArrayBG->GetDynamicArray()->HasDirtyBindings()) {
+ DAWN_TRY(UpdateDynamicArrayBindings(device, recordingContext,
+ dynamicArrayBG->GetDynamicArray()));
+ }
+ }
+
// Separate barriers with vertex stages in destination stages from all other barriers.
// This avoids creating unnecessary fragment->vertex dependencies when merging barriers.
// Eg. merging a compute->vertex barrier and a fragment->fragment barrier would create
@@ -639,7 +697,7 @@
auto PrepareResourcesForRenderPass = [](Device* device,
CommandRecordingContext* recordingContext,
const RenderPassResourceUsage& usages) -> MaybeError {
- DAWN_TRY(TransitionAndClearForSyncScope(device, recordingContext, usages));
+ DAWN_TRY(PrepareResourcesForSyncScope(device, recordingContext, usages));
// Reset all query set used on current render pass together before beginning render pass
// because the reset command must be called outside render pass
@@ -1070,7 +1128,7 @@
case Command::Dispatch: {
DispatchCmd* dispatch = mCommands.NextCommand<DispatchCmd>();
- DAWN_TRY(TransitionAndClearForSyncScope(
+ DAWN_TRY(PrepareResourcesForSyncScope(
device, recordingContext, resourceUsages.dispatchUsages[currentDispatch]));
descriptorSets.Apply(device, recordingContext, VK_PIPELINE_BIND_POINT_COMPUTE);
immediates.Apply(device, commands);
@@ -1083,7 +1141,7 @@
DispatchIndirectCmd* dispatch = mCommands.NextCommand<DispatchIndirectCmd>();
VkBuffer indirectBuffer = ToBackend(dispatch->indirectBuffer)->GetHandle();
- DAWN_TRY(TransitionAndClearForSyncScope(
+ DAWN_TRY(PrepareResourcesForSyncScope(
device, recordingContext, resourceUsages.dispatchUsages[currentDispatch]));
descriptorSets.Apply(device, recordingContext, VK_PIPELINE_BIND_POINT_COMPUTE);
immediates.Apply(device, commands);
diff --git a/src/dawn/tests/end2end/BindingArrayTests.cpp b/src/dawn/tests/end2end/BindingArrayTests.cpp
index 156cb00..fa71a05 100644
--- a/src/dawn/tests/end2end/BindingArrayTests.cpp
+++ b/src/dawn/tests/end2end/BindingArrayTests.cpp
@@ -25,6 +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 <string>
#include <vector>
#include "dawn/tests/DawnTest.h"
@@ -689,6 +690,74 @@
return device.CreateBindGroup(&descriptor);
}
+
+ // Test that `dynamicArray` (with layout `bgl` and `dynamicArrayStart`), has bindings of
+ // `wgslType` in the `expected` slots.
+ void TestHasBinding(wgpu::BindGroupLayout bgl,
+ wgpu::BindGroup dynamicArray,
+ std::vector<bool> expected,
+ uint32_t dynamicArrayStart = 0,
+ std::string wgslType = "texture_2d<f32>") {
+ // Create the test pipeline.
+ std::array<wgpu::BindGroupLayout, 2> bgls = {
+ bgl,
+ utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}}),
+ };
+ wgpu::PipelineLayoutDescriptor plDesc = {
+ .bindGroupLayoutCount = 2,
+ .bindGroupLayouts = bgls.data(),
+ };
+
+ wgpu::ShaderModule module =
+ utils::CreateShaderModule(device, R"(
+ enable chromium_experimental_dynamic_binding;
+ @group(0) @binding()" + std::to_string(dynamicArrayStart) +
+ R"() var bindings : resource_binding;
+ @group(1) @binding(0) var<storage, read_write> results : array<u32>;
+
+ @compute @workgroup_size(1) fn main() {
+ for (var i = 0u; i < arrayLength(bindings); i++) {
+ results[i] = u32(hasBinding<)" +
+ wgslType + R"(>(bindings, i));
+ }
+ }
+ )");
+
+ wgpu::ComputePipelineDescriptor csDesc = {.layout = device.CreatePipelineLayout(&plDesc),
+ .compute = {
+ .module = module,
+ }};
+ wgpu::ComputePipeline testPipeline = device.CreateComputePipeline(&csDesc);
+
+ // Create the result buffer.
+ wgpu::BufferDescriptor bDesc = {
+ .usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc,
+ .size = sizeof(uint32_t) * expected.size(),
+ };
+ wgpu::Buffer resultBuffer = device.CreateBuffer(&bDesc);
+ wgpu::BindGroup resultBG = utils::MakeBindGroup(device, bgls[1], {{0, resultBuffer}});
+
+ // Run the test.
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetBindGroup(0, dynamicArray);
+ pass.SetBindGroup(1, resultBG);
+ pass.SetPipeline(testPipeline);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+
+ wgpu::CommandBuffer commands = encoder.Finish();
+ device.GetQueue().Submit(1, &commands);
+
+ // Check we have the expected results.
+ std::vector<uint32_t> expectedU32;
+ for (bool b : expected) {
+ expectedU32.push_back(b ? 1u : 0u);
+ }
+
+ EXPECT_BUFFER_U32_RANGE_EQ(expectedU32.data(), resultBuffer, 0, expectedU32.size());
+ }
};
// Tests that creating the bind group that's only a dynamic array doesn't crash in backends.
@@ -862,19 +931,185 @@
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
+// Test WGSL `hasBinding` reflects the state of a dynamic binding array.
+TEST_P(DynamicBindingArrayTests, HasBindingOneTexturePinUnpin) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 3, {{1, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false, false, false});
+
+ // After pinning it has the one valid entry valid.
+ tex.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {false, true, false});
+
+ // After unpinning it has the no more valid entries.
+ tex.Unpin();
+ TestHasBinding(bgl, bg, {false, false, false});
+}
+
+// Test pin/unpin updating the availability takes into account the static bindings (so even if it
+// doesn't start at BindingIndex 0, things still work)
+TEST_P(DynamicBindingArrayTests, HasBindingOneTexturePinUnpinWithStaticBindings) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(
+ wgpu::DynamicBindingKind::SampledTexture, 4,
+ {{0, wgpu::ShaderStage::Compute, wgpu::TextureSampleType::UnfilterableFloat}});
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 3, {{0, tex.CreateView()}, {5, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false, false, false}, 4);
+
+ // After pinning it has the one valid entry valid.
+ tex.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {false, true, false}, 4);
+
+ // After unpinning it has the no more valid entries.
+ tex.Unpin();
+ TestHasBinding(bgl, bg, {false, false, false}, 4);
+}
+
+// Test that calling texture.Destroy() implicitly unpins it.
+TEST_P(DynamicBindingArrayTests, HasBindingOneTexturePinDestroy) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 3, {{1, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false, false, false});
+
+ // After pinning it has the one valid entry valid.
+ tex.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {false, true, false});
+
+ // After texture destruction it has the no more valid entries.
+ tex.Destroy();
+ TestHasBinding(bgl, bg, {false, false, false});
+}
+
+// Test that a texture used multiple times in the same dynamic binding array has its availability
+// correctly updated.
+TEST_P(DynamicBindingArrayTests, HasBindingSameTextureMultipleTimesPinUnpin) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 4, {{1, tex.CreateView()}, {3, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false, false, false, false});
+
+ // After pinning it has valid entries.
+ tex.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {false, true, false, true});
+
+ // After unpinning it has the no more valid entries.
+ tex.Unpin();
+ TestHasBinding(bgl, bg, {false, false, false, false});
+}
+
+// Test that creating a dynamic binding array with an already destroyed texture works, but doesn't
+// show that entry as available.
+TEST_P(DynamicBindingArrayTests, HasBindingDynamicArrayCreatedWithTextureAlreadyDestroyed) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+ tex.Destroy();
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 1, {{0, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false});
+}
+
+// Test that a texture used multiple times in the same dynamic binding array has its
+// availability correctly updated.
+TEST_P(DynamicBindingArrayTests, HasBindingSameTextureMultipleDynamicArrays) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg1 = MakeBindGroup(bgl, 3, {{1, tex.CreateView()}});
+ wgpu::BindGroup bg2 = MakeBindGroup(bgl, 1, {{0, tex.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg1, {false, false, false});
+ TestHasBinding(bgl, bg2, {false});
+
+ // After pinning it has valid entries.
+ tex.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg1, {false, true, false});
+ TestHasBinding(bgl, bg2, {true});
+
+ // After destroying on dynamic binding array, the other still has the texture available.
+ bg1.Destroy();
+ TestHasBinding(bgl, bg2, {true});
+}
+
+// Test that texture availabililty is controlled per-texture.
+TEST_P(DynamicBindingArrayTests, HasBindingMultipleTexturesInDynamicArray) {
+ wgpu::TextureDescriptor tDesc{
+ .usage = wgpu::TextureUsage::TextureBinding,
+ .size = {1, 1},
+ .format = wgpu::TextureFormat::R32Float,
+ };
+ wgpu::Texture tex0 = device.CreateTexture(&tDesc);
+ wgpu::Texture tex1 = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroupLayout bgl = MakeBindGroupLayout(wgpu::DynamicBindingKind::SampledTexture);
+ wgpu::BindGroup bg = MakeBindGroup(bgl, 2, {{0, tex0.CreateView()}, {1, tex1.CreateView()}});
+
+ // Before pinning, the bind group has no valid entries.
+ TestHasBinding(bgl, bg, {false, false});
+
+ // After pinning tex0 it has one valid entry.
+ tex0.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {true, false});
+
+ // After pinning tex1 it has two valid entries.
+ tex1.Pin(wgpu::TextureUsage::TextureBinding);
+ TestHasBinding(bgl, bg, {true, true});
+
+ // After unpinning tex0 it has only one valid entry.
+ tex0.Unpin();
+ TestHasBinding(bgl, bg, {false, true});
+}
+
+// TODO(https://crbug.com/435317394): When wgpu::BindGroup::Update() or equivalent is added, test
+// that availability is updated when entries in the dynamic binding array are updated.
+// TODO(https://crbug.com/435317394): Add tests that hasBinding() works as expected for all support
+// types in WGSL.
DAWN_INSTANTIATE_TEST(DynamicBindingArrayTests, D3D12Backend(), MetalBackend(), VulkanBackend());
diff --git a/src/tint/api/common/resource_type.h b/src/tint/api/common/resource_type.h
index 96fde44..0d23267 100644
--- a/src/tint/api/common/resource_type.h
+++ b/src/tint/api/common/resource_type.h
@@ -32,7 +32,7 @@
namespace tint {
-enum class ResourceType {
+enum class ResourceType : uint32_t {
kEmpty,
kTexture1d_f32,
