Fix Synchronization Bug For Shared Buffers in Bind Groups
Changes how shared buffers are recognized for usage to include shared
buffers used within bind groups. Includes a test.
Bug: 42241325
Change-Id: Ida5f765d596384fc0bb70d91ebd2e8260969a725
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/232256
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Brandon Jones <bajones@chromium.org>
Commit-Queue: Brandon1 Jones <brandon1.jones@intel.com>
diff --git a/src/dawn/native/d3d12/BufferD3D12.cpp b/src/dawn/native/d3d12/BufferD3D12.cpp
index 3e454e9..ab9b5f8 100644
--- a/src/dawn/native/d3d12/BufferD3D12.cpp
+++ b/src/dawn/native/d3d12/BufferD3D12.cpp
@@ -340,6 +340,10 @@
bool Buffer::TrackUsageAndGetResourceBarrier(CommandRecordingContext* commandContext,
D3D12_RESOURCE_BARRIER* barrier,
wgpu::BufferUsage newUsage) {
+ if (mResourceAllocation.GetInfo().mMethod == AllocationMethod::kExternal) {
+ commandContext->AddToSharedBufferList(this);
+ }
+
// Track the underlying heap to ensure residency.
// There may be no heap if the allocation is an external one.
Heap* heap = ToBackend(mResourceAllocation.GetResourceHeap());
@@ -424,10 +428,6 @@
void Buffer::TrackUsageAndTransitionNow(CommandRecordingContext* commandContext,
wgpu::BufferUsage newUsage) {
- if (mResourceAllocation.GetInfo().mMethod == AllocationMethod::kExternal) {
- commandContext->AddToSharedBufferList(this);
- }
-
D3D12_RESOURCE_BARRIER barrier;
if (TrackUsageAndGetResourceBarrier(commandContext, &barrier, newUsage)) {
diff --git a/src/dawn/tests/white_box/SharedBufferMemoryTests.cpp b/src/dawn/tests/white_box/SharedBufferMemoryTests.cpp
index 5b01185..f9d41f9 100644
--- a/src/dawn/tests/white_box/SharedBufferMemoryTests.cpp
+++ b/src/dawn/tests/white_box/SharedBufferMemoryTests.cpp
@@ -573,6 +573,71 @@
}
}
+// Test to ensure that using a shared buffer in a bind group will trigger a wait for the fence
+// provided to BeginAccess.
+TEST_P(SharedBufferMemoryTests, UseInPassEnsureSynchronization) {
+ wgpu::SharedBufferMemory memory =
+ GetParam().mBackend->CreateSharedBufferMemory(device, kStorageUsages, kBufferSize);
+ wgpu::Buffer buffer = memory.CreateBuffer();
+
+ wgpu::SharedBufferMemoryBeginAccessDescriptor beginAccessDesc;
+ beginAccessDesc.initialized = true;
+ memory.BeginAccess(buffer, &beginAccessDesc);
+
+ wgpu::Buffer srcBuffer =
+ utils::CreateBufferFromData(device, &kBufferData, kBufferSize, wgpu::BufferUsage::CopySrc);
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ encoder.CopyBufferToBuffer(srcBuffer, 0, buffer, 0, kBufferSize);
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ wgpu::SharedBufferMemoryEndAccessState endState;
+ memory.EndAccess(buffer, &endState);
+
+ // Pass fences from the previous operation to the next BeginAccessDescriptor to ensure
+ // operations are complete.
+ std::vector<wgpu::SharedFence> sharedFences(endState.fenceCount);
+ for (size_t j = 0; j < endState.fenceCount; ++j) {
+ sharedFences[j] = GetParam().mBackend->ImportFenceTo(device, endState.fences[j]);
+ }
+ beginAccessDesc.fenceCount = sharedFences.size();
+ beginAccessDesc.fences = sharedFences.data();
+ beginAccessDesc.signaledValues = endState.signaledValues;
+ beginAccessDesc.initialized = true;
+ wgpu::Buffer buffer2 = memory.CreateBuffer();
+ memory.BeginAccess(buffer2, &beginAccessDesc);
+
+ wgpu::ComputePipelineDescriptor pipelineDescriptor;
+
+ // This compute shader reads from the shared storage buffer and increments it by one.
+ pipelineDescriptor.compute.module = utils::CreateShaderModule(device, R"(
+ struct OutputBuffer {
+ value : u32
+ }
+
+ @group(0) @binding(0) var<storage, read_write> outputBuffer : OutputBuffer;
+
+ @compute @workgroup_size(1) fn main() {
+ outputBuffer.value = outputBuffer.value + 1u;
+ })");
+
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor);
+ wgpu::BindGroup bindGroup =
+ utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer2}});
+ encoder = device.CreateCommandEncoder();
+ wgpu::CommandBuffer commands;
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+ commands = encoder.Finish();
+ queue.Submit(1, &commands);
+
+ // The storage buffer should have been incremented by one in the compute shader.
+ EXPECT_BUFFER_U32_EQ(kBufferData + 1, buffer2, 0);
+}
+
// Test to ensure WriteBuffer waits on a fence provided to BeginAccess.
TEST_P(SharedBufferMemoryTests, WriteBufferEnsureSynchronization) {
wgpu::SharedBufferMemory memory =
