From c872e6a5520f3fd3f33bb2ed04a8703755b87301 Mon Sep 17 00:00:00 2001 From: Peng Huang Date: Wed, 10 May 2023 17:21:53 +0000 Subject: [PATCH] d3d11: support uniform buffer with other GPU accelerated usage d3d11 doesn't allow creating a constant buffer with other GPU accelerated usage. This CL workarounds problem by creating two buffers one for uniform buffer usage, one for other usage, and copy content to uniform buffer when it is needed. Bug: dawn:1755 Bug: dawn:1798 Bug: dawn:1721 Change-Id: I26bfee1cca2204f6464ba611872c490165e97f68 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/132020 Commit-Queue: Peng Huang Reviewed-by: Austin Eng Kokoro: Kokoro --- .../native/d3d11/BindGroupTrackerD3D11.cpp | 5 +- src/dawn/native/d3d11/BufferD3D11.cpp | 274 +++++++++------ src/dawn/native/d3d11/BufferD3D11.h | 24 +- src/dawn/native/d3d11/CommandBufferD3D11.cpp | 10 +- .../d3d11/CommandRecordingContextD3D11.cpp | 2 +- src/dawn/native/d3d11/TextureD3D11.h | 1 - src/dawn/tests/BUILD.gn | 4 + src/dawn/tests/white_box/D3D11BufferTests.cpp | 312 ++++++++++++++++++ 8 files changed, 512 insertions(+), 120 deletions(-) create mode 100644 src/dawn/tests/white_box/D3D11BufferTests.cpp diff --git a/src/dawn/native/d3d11/BindGroupTrackerD3D11.cpp b/src/dawn/native/d3d11/BindGroupTrackerD3D11.cpp index b7628cdbe7..8fd2d60b95 100644 --- a/src/dawn/native/d3d11/BindGroupTrackerD3D11.cpp +++ b/src/dawn/native/d3d11/BindGroupTrackerD3D11.cpp @@ -132,7 +132,6 @@ MaybeError BindGroupTracker::ApplyBindGroup(BindGroupIndex index) { switch (bindingInfo.bindingType) { case BindingInfoType::Buffer: { BufferBinding binding = group->GetBindingAsBufferBinding(bindingIndex); - ID3D11Buffer* d3d11Buffer = ToBackend(binding.buffer)->GetD3D11Buffer(); auto offset = binding.offset; if (bindingInfo.buffer.hasDynamicOffset) { // Dynamic buffers are packed at the front of BindingIndices. @@ -141,6 +140,9 @@ MaybeError BindGroupTracker::ApplyBindGroup(BindGroupIndex index) { switch (bindingInfo.buffer.type) { case wgpu::BufferBindingType::Uniform: { + ToBackend(binding.buffer)->EnsureConstantBufferIsUpdated(mCommandContext); + ID3D11Buffer* d3d11Buffer = + ToBackend(binding.buffer)->GetD3D11ConstantBuffer(); // https://learn.microsoft.com/en-us/windows/win32/api/d3d11_1/nf-d3d11_1-id3d11devicecontext1-vssetconstantbuffers1 // Offset and size are measured in shader constants, which are 16 bytes // (4*32-bit components). And the offsets and counts must be multiples @@ -175,6 +177,7 @@ MaybeError BindGroupTracker::ApplyBindGroup(BindGroupIndex index) { DAWN_TRY_ASSIGN( d3d11UAV, ToBackend(binding.buffer) ->CreateD3D11UnorderedAccessView1(offset, binding.size)); + ToBackend(binding.buffer)->MarkMutated(); if (bindingInfo.visibility & wgpu::ShaderStage::Fragment) { deviceContext1->OMSetRenderTargetsAndUnorderedAccessViews( D3D11_KEEP_RENDER_TARGETS_AND_DEPTH_STENCIL, nullptr, nullptr, diff --git a/src/dawn/native/d3d11/BufferD3D11.cpp b/src/dawn/native/d3d11/BufferD3D11.cpp index b399870bef..233c195e0d 100644 --- a/src/dawn/native/d3d11/BufferD3D11.cpp +++ b/src/dawn/native/d3d11/BufferD3D11.cpp @@ -33,19 +33,8 @@ namespace dawn::native::d3d11 { namespace { -MaybeError ValidationUsage(wgpu::BufferUsage usage) { - // https://learn.microsoft.com/en-us/windows/win32/api/d3d11/ne-d3d11-d3d11_bind_flag - // D3D11 doesn't support constants buffers with other accelerated GPU usages. - // TODO(dawn:1755): find a way to workaround this D3D11 limitation. - constexpr wgpu::BufferUsage kAllowedUniformBufferUsages = - wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Uniform; - - DAWN_INVALID_IF( - usage & wgpu::BufferUsage::Uniform && !IsSubset(usage, kAllowedUniformBufferUsages), - "Buffer usage can't be both uniform and other accelerated usages with D3D11"); - - return {}; -} +constexpr wgpu::BufferUsage kD3D11AllowedUniformBufferUsages = + wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; // Resource usage Default Dynamic Immutable Staging // ------------------------------------------------------------ @@ -74,19 +63,19 @@ UINT D3D11BufferBindFlags(wgpu::BufferUsage usage) { UINT bindFlags = 0; if (usage & (wgpu::BufferUsage::Vertex)) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_VERTEX_BUFFER; + bindFlags |= D3D11_BIND_VERTEX_BUFFER; } if (usage & wgpu::BufferUsage::Index) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_INDEX_BUFFER; + bindFlags |= D3D11_BIND_INDEX_BUFFER; } if (usage & (wgpu::BufferUsage::Uniform)) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_CONSTANT_BUFFER; + bindFlags |= D3D11_BIND_CONSTANT_BUFFER; } if (usage & (wgpu::BufferUsage::Storage | kInternalStorageBuffer)) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_UNORDERED_ACCESS; + bindFlags |= D3D11_BIND_UNORDERED_ACCESS; } if (usage & kReadOnlyStorageBuffer) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_SHADER_RESOURCE; + bindFlags |= D3D11_BIND_SHADER_RESOURCE; } constexpr wgpu::BufferUsage kCopyUsages = @@ -96,7 +85,7 @@ UINT D3D11BufferBindFlags(wgpu::BufferUsage usage) { // to copy data between buffer and texture. So the buffer needs to be bound as unordered access // view. if (IsSubset(usage, kCopyUsages)) { - bindFlags |= D3D11_BIND_FLAG::D3D11_BIND_UNORDERED_ACCESS; + bindFlags |= D3D11_BIND_UNORDERED_ACCESS; } return bindFlags; @@ -112,8 +101,7 @@ UINT D3D11CpuAccessFlags(wgpu::BufferUsage usage) { // - For texture to buffer copy, we may need copy texture to a staging (mappable) // texture, and then memcpy the data from the staging texture to the staging buffer. So // D3D11_CPU_ACCESS_WRITE is needed to MapRead usage. - cpuAccessFlags = D3D11_CPU_ACCESS_FLAG::D3D11_CPU_ACCESS_READ | - D3D11_CPU_ACCESS_FLAG::D3D11_CPU_ACCESS_WRITE; + cpuAccessFlags = D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE; } return cpuAccessFlags; } @@ -155,7 +143,6 @@ ResultOrError> Buffer::Create(Device* device, const BufferDescriptor MaybeError Buffer::Initialize(bool mappedAtCreation) { // TODO(dawn:1705): handle mappedAtCreation for NonzeroClearResourcesOnCreationForTesting - DAWN_TRY(ValidationUsage(GetUsage())); // Allocate at least 4 bytes so clamped accesses are always in bounds. uint64_t size = std::max(GetSize(), uint64_t(4u)); @@ -167,19 +154,46 @@ MaybeError Buffer::Initialize(bool mappedAtCreation) { } mAllocatedSize = Align(size, alignment); - // Create mD3d11Buffer - D3D11_BUFFER_DESC bufferDescriptor; - bufferDescriptor.ByteWidth = mAllocatedSize; - bufferDescriptor.Usage = D3D11BufferUsage(GetUsage()); - bufferDescriptor.BindFlags = D3D11BufferBindFlags(GetUsage()); - bufferDescriptor.CPUAccessFlags = D3D11CpuAccessFlags(GetUsage()); - bufferDescriptor.MiscFlags = D3D11BufferMiscFlags(GetUsage()); - bufferDescriptor.StructureByteStride = 0; + bool needsConstantBuffer = GetUsage() & wgpu::BufferUsage::Uniform; + bool onlyNeedsConstantBuffer = + needsConstantBuffer && IsSubset(GetUsage(), kD3D11AllowedUniformBufferUsages); - DAWN_TRY(CheckOutOfMemoryHRESULT(ToBackend(GetDevice()) - ->GetD3D11Device() - ->CreateBuffer(&bufferDescriptor, nullptr, &mD3d11Buffer), - "ID3D11Device::CreateBuffer")); + if (!onlyNeedsConstantBuffer) { + // Create mD3d11NonConstantBuffer + wgpu::BufferUsage nonUniformUsage = GetUsage() & ~wgpu::BufferUsage::Uniform; + D3D11_BUFFER_DESC bufferDescriptor; + bufferDescriptor.ByteWidth = mAllocatedSize; + bufferDescriptor.Usage = D3D11BufferUsage(nonUniformUsage); + bufferDescriptor.BindFlags = D3D11BufferBindFlags(nonUniformUsage); + bufferDescriptor.CPUAccessFlags = D3D11CpuAccessFlags(nonUniformUsage); + bufferDescriptor.MiscFlags = D3D11BufferMiscFlags(nonUniformUsage); + bufferDescriptor.StructureByteStride = 0; + + DAWN_TRY(CheckOutOfMemoryHRESULT( + ToBackend(GetDevice()) + ->GetD3D11Device() + ->CreateBuffer(&bufferDescriptor, nullptr, &mD3d11NonConstantBuffer), + "ID3D11Device::CreateBuffer")); + } + + if (needsConstantBuffer) { + // Create mD3d11ConstantBuffer + D3D11_BUFFER_DESC bufferDescriptor; + bufferDescriptor.ByteWidth = mAllocatedSize; + bufferDescriptor.Usage = D3D11_USAGE_DEFAULT; + bufferDescriptor.BindFlags = D3D11_BIND_CONSTANT_BUFFER; + bufferDescriptor.CPUAccessFlags = 0; + bufferDescriptor.MiscFlags = 0; + bufferDescriptor.StructureByteStride = 0; + + DAWN_TRY(CheckOutOfMemoryHRESULT( + ToBackend(GetDevice()) + ->GetD3D11Device() + ->CreateBuffer(&bufferDescriptor, nullptr, &mD3d11ConstantBuffer), + "ID3D11Device::CreateBuffer")); + } + + ASSERT(mD3d11NonConstantBuffer || mD3d11ConstantBuffer); SetLabelImpl(); return {}; @@ -201,10 +215,11 @@ MaybeError Buffer::MapInternal() { // need write permission to initialize the buffer. // TODO(dawn:1705): investigate the performance impact of mapping with D3D11_MAP_READ_WRITE. D3D11_MAPPED_SUBRESOURCE mappedResource; - DAWN_TRY(CheckHRESULT(commandContext->GetD3D11DeviceContext()->Map( - mD3d11Buffer.Get(), /*Subresource=*/0, D3D11_MAP_READ_WRITE, - /*MapFlags=*/0, &mappedResource), - "ID3D11DeviceContext::Map")); + DAWN_TRY(CheckHRESULT( + commandContext->GetD3D11DeviceContext()->Map(mD3d11NonConstantBuffer.Get(), + /*Subresource=*/0, D3D11_MAP_READ_WRITE, + /*MapFlags=*/0, &mappedResource), + "ID3D11DeviceContext::Map")); mMappedData = reinterpret_cast(mappedResource.pData); return {}; @@ -214,7 +229,8 @@ void Buffer::UnmapInternal() { DAWN_ASSERT(mMappedData); CommandRecordingContext* commandContext = ToBackend(GetDevice())->GetPendingCommandContext(); - commandContext->GetD3D11DeviceContext()->Unmap(mD3d11Buffer.Get(), /*Subresource=*/0); + commandContext->GetD3D11DeviceContext()->Unmap(mD3d11NonConstantBuffer.Get(), + /*Subresource=*/0); mMappedData = nullptr; } @@ -224,7 +240,7 @@ MaybeError Buffer::MapAtCreationImpl() { } MaybeError Buffer::MapAsyncImpl(wgpu::MapMode mode, size_t offset, size_t size) { - DAWN_ASSERT(mD3d11Buffer); + DAWN_ASSERT(mD3d11NonConstantBuffer); // TODO(dawn:1705): make sure the map call is not blocked by the GPU operations. DAWN_TRY(MapInternal()); @@ -236,7 +252,7 @@ MaybeError Buffer::MapAsyncImpl(wgpu::MapMode mode, size_t offset, size_t size) } void Buffer::UnmapImpl() { - DAWN_ASSERT(mD3d11Buffer); + DAWN_ASSERT(mD3d11NonConstantBuffer); DAWN_ASSERT(mMappedData); UnmapInternal(); } @@ -252,11 +268,13 @@ void Buffer::DestroyImpl() { if (mMappedData) { UnmapInternal(); } - mD3d11Buffer = nullptr; + mD3d11NonConstantBuffer = nullptr; } void Buffer::SetLabelImpl() { - SetDebugName(ToBackend(GetDevice()), mD3d11Buffer.Get(), "Dawn_Buffer", GetLabel()); + SetDebugName(ToBackend(GetDevice()), mD3d11NonConstantBuffer.Get(), "Dawn_Buffer", GetLabel()); + SetDebugName(ToBackend(GetDevice()), mD3d11ConstantBuffer.Get(), "Dawn_ConstantBuffer", + GetLabel()); } MaybeError Buffer::EnsureDataInitialized(CommandRecordingContext* commandContext) { @@ -309,6 +327,22 @@ MaybeError Buffer::InitializeToZero(CommandRecordingContext* commandContext) { return {}; } +void Buffer::MarkMutated() { + mConstantBufferIsUpdated = false; +} + +void Buffer::EnsureConstantBufferIsUpdated(CommandRecordingContext* commandContext) { + if (mConstantBufferIsUpdated) { + return; + } + + DAWN_ASSERT(mD3d11NonConstantBuffer); + DAWN_ASSERT(mD3d11ConstantBuffer); + commandContext->GetD3D11DeviceContext1()->CopyResource(mD3d11ConstantBuffer.Get(), + mD3d11NonConstantBuffer.Get()); + mConstantBufferIsUpdated = true; +} + ResultOrError> Buffer::CreateD3D11ShaderResourceView( uint64_t offset, uint64_t size) const { @@ -324,10 +358,11 @@ ResultOrError> Buffer::CreateD3D11ShaderResourc desc.BufferEx.NumElements = numElements; desc.BufferEx.Flags = D3D11_BUFFEREX_SRV_FLAG_RAW; ComPtr srv; - DAWN_TRY(CheckHRESULT(ToBackend(GetDevice()) - ->GetD3D11Device() - ->CreateShaderResourceView(mD3d11Buffer.Get(), &desc, &srv), - "ShaderResourceView creation")); + DAWN_TRY( + CheckHRESULT(ToBackend(GetDevice()) + ->GetD3D11Device() + ->CreateShaderResourceView(mD3d11NonConstantBuffer.Get(), &desc, &srv), + "ShaderResourceView creation")); return srv; } @@ -349,11 +384,11 @@ ResultOrError> Buffer::CreateD3D11UnorderedAc desc.Buffer.Flags = D3D11_BUFFER_UAV_FLAG_RAW; ComPtr uav; - DAWN_TRY(CheckHRESULT(ToBackend(GetDevice()) - ->GetD3D11Device5() - ->CreateUnorderedAccessView1(mD3d11Buffer.Get(), &desc, &uav), - "UnorderedAccessView creation")); - + DAWN_TRY( + CheckHRESULT(ToBackend(GetDevice()) + ->GetD3D11Device5() + ->CreateUnorderedAccessView1(mD3d11NonConstantBuffer.Get(), &desc, &uav), + "UnorderedAccessView creation")); return uav; } @@ -388,6 +423,8 @@ MaybeError Buffer::ClearInternal(CommandRecordingContext* commandContext, if (mMappedData) { memset(mMappedData + offset, clearValue, size); + // The WebGPU uniform buffer is not mappable. + ASSERT(!mD3d11ConstantBuffer); return {}; } @@ -428,6 +465,8 @@ MaybeError Buffer::WriteInternal(CommandRecordingContext* commandContext, if (scopedMap.GetMappedData()) { memcpy(scopedMap.GetMappedData() + offset, data, size); + // The WebGPU uniform buffer is not mappable. + ASSERT(!mD3d11ConstantBuffer); return {}; } @@ -436,65 +475,55 @@ MaybeError Buffer::WriteInternal(CommandRecordingContext* commandContext, ID3D11DeviceContext1* d3d11DeviceContext1 = commandContext->GetD3D11DeviceContext1(); - // For updating the full buffer, just pass nullptr as the pDstBox. - if (offset == 0 && size == GetAllocatedSize()) { - d3d11DeviceContext1->UpdateSubresource(GetD3D11Buffer(), /*DstSubresource=*/0, - /*pDstBox=*/nullptr, data, + if (mD3d11NonConstantBuffer) { + D3D11_BOX box; + box.left = offset; + box.right = offset + size; + box.top = 0; + box.bottom = 1; + box.front = 0; + box.back = 1; + d3d11DeviceContext1->UpdateSubresource(mD3d11NonConstantBuffer.Get(), /*DstSubresource=*/0, + &box, data, /*SrcRowPitch=*/0, /*SrcDepthPitch*/ 0); + if (!mD3d11ConstantBuffer) { + return {}; + } + + // if mConstantBufferIsUpdated is false, the content of mD3d11ConstantBuffer will be + // updated by EnsureConstantBufferIsUpdated() when the constant buffer is about to be used. + if (!mConstantBufferIsUpdated) { + return {}; + } + + // Copy the modified part of the mD3d11NonConstantBuffer to mD3d11ConstantBuffer. + d3d11DeviceContext1->CopySubresourceRegion( + mD3d11ConstantBuffer.Get(), /*DstSubresource=*/0, /*DstX=*/offset, + /*DstY=*/0, + /*DstZ=*/0, mD3d11NonConstantBuffer.Get(), /*SrcSubresource=*/0, &box); + return {}; } - D3D11_BOX box; - box.left = offset; - box.right = offset + size; - box.top = 0; - box.bottom = 1; - box.front = 0; - box.back = 1; + ASSERT(mD3d11ConstantBuffer); - if ((GetUsage() & wgpu::BufferUsage::Uniform)) { - if (!IsAligned(box.left, 16) || !IsAligned(box.right, 16)) { - // Create a temp staging buffer to workaround the alignment issue. - BufferDescriptor descriptor; - descriptor.size = box.right - box.left; - DAWN_ASSERT(IsAligned(descriptor.size, 4)); - descriptor.usage = wgpu::BufferUsage::MapWrite | wgpu::BufferUsage::CopySrc; - descriptor.mappedAtCreation = false; - descriptor.label = "temp staging buffer"; - Ref stagingBufferBase; - DAWN_TRY_ASSIGN(stagingBufferBase, GetDevice()->CreateBuffer(&descriptor)); - Ref stagingBuffer; - stagingBuffer = ToBackend(std::move(stagingBufferBase)); - { - ScopedMap scopedMap; - DAWN_TRY_ASSIGN(scopedMap, ScopedMap::Create(stagingBuffer.Get())); - uint8_t* mappedData = scopedMap.GetMappedData(); - DAWN_ASSERT(mappedData); - memcpy(mappedData, data, size); - } - box.left = 0; - box.right = descriptor.size; - commandContext->GetD3D11DeviceContext()->CopySubresourceRegion( - GetD3D11Buffer(), /*DstSubresource=*/0, /*DstX=*/offset, - /*DstY=*/0, - /*DstZ=*/0, stagingBuffer->GetD3D11Buffer(), /*SrcSubresource=*/0, &box); - stagingBuffer = nullptr; + // If the mD3d11NonConstantBuffer is null, we have to create a staging buffer for transfer the + // data to mD3d11ConstantBuffer, since UpdateSubresource() has many restrictions. For example, + // the size of the data has to be a multiple of 16, etc + BufferDescriptor descriptor; + descriptor.size = size; + DAWN_ASSERT(IsAligned(descriptor.size, 4)); + descriptor.usage = wgpu::BufferUsage::MapWrite | wgpu::BufferUsage::CopySrc; + descriptor.mappedAtCreation = false; + descriptor.label = "DawnWriteStagingBuffer"; + Ref stagingBuffer; + DAWN_TRY_ASSIGN(stagingBuffer, GetDevice()->CreateBuffer(&descriptor)); - } else { - // TODO(dawn:1739): check whether driver supports partial update of uniform buffer. - d3d11DeviceContext1->UpdateSubresource1(GetD3D11Buffer(), /*DstSubresource=*/0, &box, - data, - /*SrcRowPitch=*/0, - /*SrcDepthPitch*/ 0, D3D11_COPY_NO_OVERWRITE); - } - } else { - d3d11DeviceContext1->UpdateSubresource(GetD3D11Buffer(), /*DstSubresource=*/0, &box, data, - /*SrcRowPitch=*/0, - /*SrcDepthPitch*/ 0); - } + DAWN_TRY(ToBackend(stagingBuffer)->WriteInternal(commandContext, 0, data, size)); - return {}; + return Buffer::CopyInternal(commandContext, ToBackend(stagingBuffer.Get()), /*sourceOffset=*/0, + /*size=*/size, this, offset); } // static @@ -509,7 +538,16 @@ MaybeError Buffer::Copy(CommandRecordingContext* commandContext, DAWN_TRY(source->EnsureDataInitialized(commandContext)); DAWN_TRY( destination->EnsureDataInitializedAsDestination(commandContext, destinationOffset, size)); + return CopyInternal(commandContext, source, sourceOffset, size, destination, destinationOffset); +} +// static +MaybeError Buffer::CopyInternal(CommandRecordingContext* commandContext, + Buffer* source, + uint64_t sourceOffset, + size_t size, + Buffer* destination, + uint64_t destinationOffset) { D3D11_BOX srcBox; srcBox.left = sourceOffset; srcBox.right = sourceOffset + size; @@ -517,10 +555,32 @@ MaybeError Buffer::Copy(CommandRecordingContext* commandContext, srcBox.bottom = 1; srcBox.front = 0; srcBox.back = 1; - commandContext->GetD3D11DeviceContext()->CopySubresourceRegion( - destination->mD3d11Buffer.Get(), /*DstSubresource=*/0, /*DstX=*/destinationOffset, - /*DstY=*/0, - /*DstZ=*/0, source->mD3d11Buffer.Get(), /*SrcSubresource=*/0, &srcBox); + ID3D11Buffer* d3d11SourceBuffer = source->mD3d11NonConstantBuffer + ? source->mD3d11NonConstantBuffer.Get() + : source->mD3d11ConstantBuffer.Get(); + ASSERT(d3d11SourceBuffer); + + if (destination->mD3d11NonConstantBuffer) { + commandContext->GetD3D11DeviceContext()->CopySubresourceRegion( + destination->mD3d11NonConstantBuffer.Get(), /*DstSubresource=*/0, + /*DstX=*/destinationOffset, + /*DstY=*/0, + /*DstZ=*/0, d3d11SourceBuffer, /*SrcSubresource=*/0, &srcBox); + } + + // if mConstantBufferIsUpdated is false, the content of mD3d11ConstantBuffer will be + // updated by EnsureConstantBufferIsUpdated() when the constant buffer is about to be used. + if (!destination->mConstantBufferIsUpdated) { + return {}; + } + + if (destination->mD3d11ConstantBuffer) { + commandContext->GetD3D11DeviceContext()->CopySubresourceRegion( + destination->mD3d11ConstantBuffer.Get(), /*DstSubresource=*/0, + /*DstX=*/destinationOffset, + /*DstY=*/0, + /*DstZ=*/0, d3d11SourceBuffer, /*SrcSubresource=*/0, &srcBox); + } return {}; } diff --git a/src/dawn/native/d3d11/BufferD3D11.h b/src/dawn/native/d3d11/BufferD3D11.h index d8633179cb..aac288d817 100644 --- a/src/dawn/native/d3d11/BufferD3D11.h +++ b/src/dawn/native/d3d11/BufferD3D11.h @@ -40,14 +40,19 @@ class Buffer final : public BufferBase { // Dawn API void SetLabelImpl() override; - ID3D11Buffer* GetD3D11Buffer() const { return mD3d11Buffer.Get(); } + ID3D11Buffer* GetD3D11ConstantBuffer() const { return mD3d11ConstantBuffer.Get(); } + ID3D11Buffer* GetD3D11NonConstantBuffer() const { return mD3d11NonConstantBuffer.Get(); } + // Mark the mD3d11NonConstantBuffer is mutated by shaders, if mD3d11ConstantBuffer exists, + // it will be synced with mD3d11NonConstantBuffer before binding it to the constant buffer slot. + void MarkMutated(); + // Update content of the mD3d11ConstantBuffer from mD3d11NonConstantBuffer if needed. + void EnsureConstantBufferIsUpdated(CommandRecordingContext* commandContext); ResultOrError> CreateD3D11ShaderResourceView( uint64_t offset, uint64_t size) const; ResultOrError> CreateD3D11UnorderedAccessView1( uint64_t offset, uint64_t size) const; - MaybeError Clear(CommandRecordingContext* commandContext, uint8_t clearValue, uint64_t offset, @@ -116,9 +121,18 @@ class Buffer final : public BufferBase { uint64_t bufferOffset, const void* data, size_t size); - - // The buffer object can be used as vertex, index, uniform, storage, or indirect buffer. - ComPtr mD3d11Buffer; + // Copy the buffer without checking if the buffer is initialized. + static MaybeError CopyInternal(CommandRecordingContext* commandContext, + Buffer* source, + uint64_t sourceOffset, + size_t size, + Buffer* destination, + uint64_t destinationOffset); + // The buffer object for constant buffer usage. + ComPtr mD3d11ConstantBuffer; + // The buffer object for non-constant buffer usages(e.g. storage buffer, vertex buffer, etc.) + ComPtr mD3d11NonConstantBuffer; + bool mConstantBufferIsUpdated = true; uint8_t* mMappedData = nullptr; }; diff --git a/src/dawn/native/d3d11/CommandBufferD3D11.cpp b/src/dawn/native/d3d11/CommandBufferD3D11.cpp index 1fc3296b55..919e317522 100644 --- a/src/dawn/native/d3d11/CommandBufferD3D11.cpp +++ b/src/dawn/native/d3d11/CommandBufferD3D11.cpp @@ -358,7 +358,7 @@ MaybeError CommandBuffer::ExecuteComputePass(CommandRecordingContext* commandCon } commandContext->GetD3D11DeviceContext()->DispatchIndirect( - indirectBuffer->GetD3D11Buffer(), dispatch->indirectOffset); + indirectBuffer->GetD3D11NonConstantBuffer(), dispatch->indirectOffset); break; } @@ -533,7 +533,7 @@ MaybeError CommandBuffer::ExecuteRenderPass(BeginRenderPassCmd* renderPass, } commandContext->GetD3D11DeviceContext()->DrawInstancedIndirect( - indirectBuffer->GetD3D11Buffer(), draw->indirectOffset); + indirectBuffer->GetD3D11NonConstantBuffer(), draw->indirectOffset); break; } @@ -559,7 +559,7 @@ MaybeError CommandBuffer::ExecuteRenderPass(BeginRenderPassCmd* renderPass, } commandContext->GetD3D11DeviceContext()->DrawIndexedInstancedIndirect( - indirectBuffer->GetD3D11Buffer(), draw->indirectOffset); + indirectBuffer->GetD3D11NonConstantBuffer(), draw->indirectOffset); break; } @@ -594,7 +594,7 @@ MaybeError CommandBuffer::ExecuteRenderPass(BeginRenderPassCmd* renderPass, DXGI_FORMAT indexBufferFormat = DXGIIndexFormat(cmd->format); commandContext->GetD3D11DeviceContext()->IASetIndexBuffer( - ToBackend(cmd->buffer)->GetD3D11Buffer(), indexBufferFormat, + ToBackend(cmd->buffer)->GetD3D11NonConstantBuffer(), indexBufferFormat, indexBufferBaseOffset); break; @@ -602,7 +602,7 @@ MaybeError CommandBuffer::ExecuteRenderPass(BeginRenderPassCmd* renderPass, case Command::SetVertexBuffer: { SetVertexBufferCmd* cmd = iter->NextCommand(); - ID3D11Buffer* buffer = ToBackend(cmd->buffer)->GetD3D11Buffer(); + ID3D11Buffer* buffer = ToBackend(cmd->buffer)->GetD3D11NonConstantBuffer(); vertexBufferTracker.OnSetVertexBuffer(cmd->slot, buffer, cmd->offset); break; } diff --git a/src/dawn/native/d3d11/CommandRecordingContextD3D11.cpp b/src/dawn/native/d3d11/CommandRecordingContextD3D11.cpp index bfbc052731..a7a3c76d58 100644 --- a/src/dawn/native/d3d11/CommandRecordingContextD3D11.cpp +++ b/src/dawn/native/d3d11/CommandRecordingContextD3D11.cpp @@ -63,7 +63,7 @@ MaybeError CommandRecordingContext::Open(Device* device) { // Always bind the uniform buffer to the reserved slot for all pipelines. // This buffer will be updated with the correct values before each draw or dispatch call. - ID3D11Buffer* bufferPtr = mUniformBuffer->GetD3D11Buffer(); + ID3D11Buffer* bufferPtr = mUniformBuffer->GetD3D11ConstantBuffer(); mD3D11DeviceContext4->VSSetConstantBuffers(PipelineLayout::kReservedConstantBufferSlot, 1, &bufferPtr); mD3D11DeviceContext4->CSSetConstantBuffers(PipelineLayout::kReservedConstantBufferSlot, 1, diff --git a/src/dawn/native/d3d11/TextureD3D11.h b/src/dawn/native/d3d11/TextureD3D11.h index 06e84a21ca..ecf08c44d1 100644 --- a/src/dawn/native/d3d11/TextureD3D11.h +++ b/src/dawn/native/d3d11/TextureD3D11.h @@ -61,7 +61,6 @@ class Texture final : public d3d::Texture { D3D11_DEPTH_STENCIL_VIEW_DESC GetDSVDescriptor(const SubresourceRange& range, bool depthReadOnly, bool stencilReadOnly) const; - MaybeError EnsureSubresourceContentInitialized(CommandRecordingContext* commandContext, const SubresourceRange& range); diff --git a/src/dawn/tests/BUILD.gn b/src/dawn/tests/BUILD.gn index c2735ba224..0d387472b1 100644 --- a/src/dawn/tests/BUILD.gn +++ b/src/dawn/tests/BUILD.gn @@ -664,6 +664,10 @@ source_set("white_box_tests_sources") { "white_box/QueryInternalShaderTests.cpp", ] + if (dawn_enable_d3d11) { + sources += [ "white_box/D3D11BufferTests.cpp" ] + } + if (dawn_enable_d3d12) { sources += [ "white_box/D3D12DescriptorHeapTests.cpp", diff --git a/src/dawn/tests/white_box/D3D11BufferTests.cpp b/src/dawn/tests/white_box/D3D11BufferTests.cpp new file mode 100644 index 0000000000..4ece2ef9d4 --- /dev/null +++ b/src/dawn/tests/white_box/D3D11BufferTests.cpp @@ -0,0 +1,312 @@ +// Copyright 2023 The Dawn Authors +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include + +#include "dawn/native/D3D11Backend.h" +#include "dawn/native/d3d11/BufferD3D11.h" +#include "dawn/native/d3d11/DeviceD3D11.h" +#include "dawn/tests/DawnTest.h" +#include "dawn/utils/ComboRenderPipelineDescriptor.h" +#include "dawn/utils/WGPUHelpers.h" + +class D3D11BufferTests : public DawnTest { + protected: + void SetUp() override { + DawnTest::SetUp(); + DAWN_TEST_UNSUPPORTED_IF(UsesWire()); + } + + wgpu::Buffer CreateBuffer(uint32_t bufferSize, wgpu::BufferUsage usage) { + wgpu::BufferDescriptor descriptor; + + descriptor.size = bufferSize; + descriptor.usage = usage; + + return device.CreateBuffer(&descriptor); + } + + ID3D11Device* GetD3D11Device() { + return dawn::native::d3d11::ToBackend(dawn::native::FromAPI((device.Get()))) + ->GetD3D11Device(); + } + + template + void CheckBuffer(ID3D11Buffer* buffer, std::vector expectedData, uint32_t offset = 0) { + D3D11_BUFFER_DESC bufferDesc; + buffer->GetDesc(&bufferDesc); + EXPECT_GE(bufferDesc.ByteWidth, (expectedData.size() + offset) * sizeof(T)); + + // Create D3D11 staging buffer + D3D11_BUFFER_DESC desc; + desc.ByteWidth = expectedData.size() * sizeof(T); + desc.Usage = D3D11_USAGE_STAGING; + desc.BindFlags = 0; + desc.CPUAccessFlags = D3D11_CPU_ACCESS_READ; + desc.MiscFlags = 0; + desc.StructureByteStride = 0; + + ComPtr stagingBuffer; + ASSERT_HRESULT_SUCCEEDED(GetD3D11Device()->CreateBuffer(&desc, nullptr, &stagingBuffer)); + + ID3D11DeviceContext* deviceContext; + GetD3D11Device()->GetImmediateContext(&deviceContext); + + // Copy buffer to staging buffer + D3D11_BOX srcBox; + srcBox.left = offset * sizeof(T); + srcBox.right = (offset + expectedData.size()) * sizeof(T); + srcBox.top = 0; + srcBox.bottom = 1; + srcBox.front = 0; + srcBox.back = 1; + deviceContext->CopySubresourceRegion(stagingBuffer.Get(), 0, 0, 0, 0, buffer, 0, &srcBox); + + // Map staging buffer + D3D11_MAPPED_SUBRESOURCE mappedResource; + ASSERT_HRESULT_SUCCEEDED( + deviceContext->Map(stagingBuffer.Get(), 0, D3D11_MAP_READ, 0, &mappedResource)); + + // Check data + const T* actualData = reinterpret_cast(mappedResource.pData); + for (size_t i = 0; i < expectedData.size(); ++i) { + EXPECT_EQ(expectedData[i], actualData[i]); + } + + // Unmap staging buffer + deviceContext->Unmap(stagingBuffer.Get(), 0); + } +}; + +// Test that creating a uniform buffer +TEST_P(D3D11BufferTests, CreateUniformBuffer) { + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_EQ(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = + wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_EQ(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Vertex; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Index; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Indirect; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Storage; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } + { + wgpu::BufferUsage usage = wgpu::BufferUsage::Storage; + wgpu::Buffer buffer = CreateBuffer(4, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_EQ(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + } +} + +// Test Buffer::Write() +TEST_P(D3D11BufferTests, WriteUniformBuffer) { + { + std::vector data = {0x12, 0x34, 0x56, 0x78}; + wgpu::BufferUsage usage = + wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; + wgpu::Buffer buffer = CreateBuffer(data.size(), usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_EQ(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + + queue.WriteBuffer(buffer, 0, data.data(), data.size()); + EXPECT_BUFFER_U8_RANGE_EQ(data.data(), buffer, 0, data.size()); + + CheckBuffer(d3d11Buffer->GetD3D11ConstantBuffer(), data); + } + { + std::vector data = {0x12, 0x34, 0x56, 0x78}; + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Vertex | + wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; + wgpu::Buffer buffer = CreateBuffer(data.size(), usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + + queue.WriteBuffer(buffer, 0, data.data(), data.size()); + EXPECT_BUFFER_U8_RANGE_EQ(data.data(), buffer, 0, data.size()); + + // both buffers should be updated. + CheckBuffer(d3d11Buffer->GetD3D11NonConstantBuffer(), data); + CheckBuffer(d3d11Buffer->GetD3D11ConstantBuffer(), data); + } +} + +// Test UAV write +TEST_P(D3D11BufferTests, WriteUniformBufferWithComputeShader) { + constexpr size_t kNumValues = 100; + std::vector data(kNumValues, 0x12345678); + uint64_t bufferSize = static_cast(data.size() * sizeof(uint32_t)); + wgpu::BufferUsage usage = wgpu::BufferUsage::Uniform | wgpu::BufferUsage::Storage | + wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc; + wgpu::Buffer buffer = CreateBuffer(bufferSize, usage); + dawn::native::d3d11::Buffer* d3d11Buffer = + dawn::native::d3d11::ToBackend(dawn::native::FromAPI(buffer.Get())); + + EXPECT_NE(d3d11Buffer->GetD3D11NonConstantBuffer(), nullptr); + EXPECT_NE(d3d11Buffer->GetD3D11ConstantBuffer(), nullptr); + + queue.WriteBuffer(buffer, 0, data.data(), bufferSize); + EXPECT_BUFFER_U32_RANGE_EQ(data.data(), buffer, 0, data.size()); + + CheckBuffer(d3d11Buffer->GetD3D11NonConstantBuffer(), data); + CheckBuffer(d3d11Buffer->GetD3D11ConstantBuffer(), data); + + // Fill the buffer with 0x11223344 with a compute shader + { + wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( + struct Buf { + data : array + } + + @group(0) @binding(0) var buf : Buf; + + @compute @workgroup_size(1) + fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3u) { + buf.data[GlobalInvocationID.x] = + vec4u(0x11223344u, 0x11223344u, 0x11223344u, 0x11223344u); + } + )"); + + wgpu::ComputePipelineDescriptor pipelineDesc = {}; + pipelineDesc.compute.module = module; + pipelineDesc.compute.entryPoint = "main"; + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc); + + wgpu::BindGroup bindGroupA = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + { + {0, buffer, 0, bufferSize}, + }); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroupA); + pass.DispatchWorkgroups(kNumValues / 4); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + std::vector expectedData(kNumValues, 0x11223344); + EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), buffer, 0, expectedData.size()); + // The non-constant buffer should be updated. + CheckBuffer(d3d11Buffer->GetD3D11NonConstantBuffer(), expectedData); + // The constant buffer should not be updated, until the constant buffer is used a pipeline + CheckBuffer(d3d11Buffer->GetD3D11ConstantBuffer(), data); + } + + // Copy the uniform buffer content to a new buffer with Compute shader + { + wgpu::Buffer newBuffer = + CreateBuffer(bufferSize, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc); + wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( + struct Buf { + data : array + } + + @group(0) @binding(0) var src : Buf; + @group(0) @binding(1) var dst : Buf; + + @compute @workgroup_size(1) + fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3u) { + dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x]; + } + )"); + + wgpu::ComputePipelineDescriptor pipelineDesc = {}; + pipelineDesc.compute.module = module; + pipelineDesc.compute.entryPoint = "main"; + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc); + + wgpu::BindGroup bindGroupA = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + { + {0, buffer, 0, bufferSize}, + {1, newBuffer, 0, bufferSize}, + }); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroupA); + pass.DispatchWorkgroups(kNumValues / 4); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + std::vector expectedData(kNumValues, 0x11223344); + EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), buffer, 0, expectedData.size()); + EXPECT_BUFFER_U32_RANGE_EQ(expectedData.data(), newBuffer, 0, expectedData.size()); + + // The non-constant buffer should be updated. + CheckBuffer(d3d11Buffer->GetD3D11NonConstantBuffer(), expectedData); + // The constant buffer should be updated too. + CheckBuffer(d3d11Buffer->GetD3D11ConstantBuffer(), expectedData); + } +} + +DAWN_INSTANTIATE_TEST(D3D11BufferTests, D3D11Backend());