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 <penghuang@chromium.org>
Reviewed-by: Austin Eng <enga@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>

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,

src/dawn/native/d3d11/BufferD3D11.cpp

@@ -33,19 +33,8 @@
 namespace dawn::native::d3d11 {
 namespace {
 
-MaybeError ValidationUsage(wgpu::BufferUsage usage) {
+constexpr wgpu::BufferUsage kD3D11AllowedUniformBufferUsages =
-    // https://learn.microsoft.com/en-us/windows/win32/api/d3d11/ne-d3d11-d3d11_bind_flag
+    wgpu::BufferUsage::Uniform | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc;
-    // 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 {};
-}
 
 // 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 |
+        cpuAccessFlags = D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE;
-                         D3D11_CPU_ACCESS_FLAG::D3D11_CPU_ACCESS_WRITE;
     }
     return cpuAccessFlags;
 }
@@ -155,7 +143,6 @@ ResultOrError<Ref<Buffer>> 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
+    bool needsConstantBuffer = GetUsage() & wgpu::BufferUsage::Uniform;
-    D3D11_BUFFER_DESC bufferDescriptor;
+    bool onlyNeedsConstantBuffer =
-    bufferDescriptor.ByteWidth = mAllocatedSize;
+        needsConstantBuffer && IsSubset(GetUsage(), kD3D11AllowedUniformBufferUsages);
-    bufferDescriptor.Usage = D3D11BufferUsage(GetUsage());
-    bufferDescriptor.BindFlags = D3D11BufferBindFlags(GetUsage());
-    bufferDescriptor.CPUAccessFlags = D3D11CpuAccessFlags(GetUsage());
-    bufferDescriptor.MiscFlags = D3D11BufferMiscFlags(GetUsage());
-    bufferDescriptor.StructureByteStride = 0;
 
-    DAWN_TRY(CheckOutOfMemoryHRESULT(ToBackend(GetDevice())
+    if (!onlyNeedsConstantBuffer) {
-                                         ->GetD3D11Device()
+        // Create mD3d11NonConstantBuffer
-                                         ->CreateBuffer(&bufferDescriptor, nullptr, &mD3d11Buffer),
+        wgpu::BufferUsage nonUniformUsage = GetUsage() & ~wgpu::BufferUsage::Uniform;
-                                     "ID3D11Device::CreateBuffer"));
+        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(
+    DAWN_TRY(CheckHRESULT(
-                              mD3d11Buffer.Get(), /*Subresource=*/0, D3D11_MAP_READ_WRITE,
+        commandContext->GetD3D11DeviceContext()->Map(mD3d11NonConstantBuffer.Get(),
-                              /*MapFlags=*/0, &mappedResource),
+                                                     /*Subresource=*/0, D3D11_MAP_READ_WRITE,
-                          "ID3D11DeviceContext::Map"));
+                                                     /*MapFlags=*/0, &mappedResource),
+        "ID3D11DeviceContext::Map"));
     mMappedData = reinterpret_cast<uint8_t*>(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<ComPtr<ID3D11ShaderResourceView>> Buffer::CreateD3D11ShaderResourceView(
     uint64_t offset,
     uint64_t size) const {
@@ -324,10 +358,11 @@ ResultOrError<ComPtr<ID3D11ShaderResourceView>> Buffer::CreateD3D11ShaderResourc
     desc.BufferEx.NumElements = numElements;
     desc.BufferEx.Flags = D3D11_BUFFEREX_SRV_FLAG_RAW;
     ComPtr<ID3D11ShaderResourceView> srv;
-    DAWN_TRY(CheckHRESULT(ToBackend(GetDevice())
+    DAWN_TRY(
-                              ->GetD3D11Device()
+        CheckHRESULT(ToBackend(GetDevice())
-                              ->CreateShaderResourceView(mD3d11Buffer.Get(), &desc, &srv),
+                         ->GetD3D11Device()
-                          "ShaderResourceView creation"));
+                         ->CreateShaderResourceView(mD3d11NonConstantBuffer.Get(), &desc, &srv),
+                     "ShaderResourceView creation"));
 
     return srv;
 }
@@ -349,11 +384,11 @@ ResultOrError<ComPtr<ID3D11UnorderedAccessView1>> Buffer::CreateD3D11UnorderedAc
     desc.Buffer.Flags = D3D11_BUFFER_UAV_FLAG_RAW;
 
     ComPtr<ID3D11UnorderedAccessView1> uav;
-    DAWN_TRY(CheckHRESULT(ToBackend(GetDevice())
+    DAWN_TRY(
-                              ->GetD3D11Device5()
+        CheckHRESULT(ToBackend(GetDevice())
-                              ->CreateUnorderedAccessView1(mD3d11Buffer.Get(), &desc, &uav),
+                         ->GetD3D11Device5()
-                          "UnorderedAccessView creation"));
+                         ->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 (mD3d11NonConstantBuffer) {
-    if (offset == 0 && size == GetAllocatedSize()) {
+        D3D11_BOX box;
-        d3d11DeviceContext1->UpdateSubresource(GetD3D11Buffer(), /*DstSubresource=*/0,
+        box.left = offset;
-                                               /*pDstBox=*/nullptr, data,
+        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;
+    ASSERT(mD3d11ConstantBuffer);
-    box.left = offset;
-    box.right = offset + size;
-    box.top = 0;
-    box.bottom = 1;
-    box.front = 0;
-    box.back = 1;
 
-    if ((GetUsage() & wgpu::BufferUsage::Uniform)) {
+    // If the mD3d11NonConstantBuffer is null, we have to create a staging buffer for transfer the
-        if (!IsAligned(box.left, 16) || !IsAligned(box.right, 16)) {
+    // data to mD3d11ConstantBuffer, since UpdateSubresource() has many restrictions. For example,
-            // Create a temp staging buffer to workaround the alignment issue.
+    // the size of the data has to be a multiple of 16, etc
-            BufferDescriptor descriptor;
+    BufferDescriptor descriptor;
-            descriptor.size = box.right - box.left;
+    descriptor.size = size;
-            DAWN_ASSERT(IsAligned(descriptor.size, 4));
+    DAWN_ASSERT(IsAligned(descriptor.size, 4));
-            descriptor.usage = wgpu::BufferUsage::MapWrite | wgpu::BufferUsage::CopySrc;
+    descriptor.usage = wgpu::BufferUsage::MapWrite | wgpu::BufferUsage::CopySrc;
-            descriptor.mappedAtCreation = false;
+    descriptor.mappedAtCreation = false;
-            descriptor.label = "temp staging buffer";
+    descriptor.label = "DawnWriteStagingBuffer";
-            Ref<BufferBase> stagingBufferBase;
+    Ref<BufferBase> stagingBuffer;
-            DAWN_TRY_ASSIGN(stagingBufferBase, GetDevice()->CreateBuffer(&descriptor));
+    DAWN_TRY_ASSIGN(stagingBuffer, GetDevice()->CreateBuffer(&descriptor));
-            Ref<Buffer> 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;
 
-        } else {
+    DAWN_TRY(ToBackend(stagingBuffer)->WriteInternal(commandContext, 0, data, size));
-            // 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);
-    }
 
-    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(
+    ID3D11Buffer* d3d11SourceBuffer = source->mD3d11NonConstantBuffer
-        destination->mD3d11Buffer.Get(), /*DstSubresource=*/0, /*DstX=*/destinationOffset,
+                                          ? source->mD3d11NonConstantBuffer.Get()
-        /*DstY=*/0,
+                                          : source->mD3d11ConstantBuffer.Get();
-        /*DstZ=*/0, source->mD3d11Buffer.Get(), /*SrcSubresource=*/0, &srcBox);
+    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 {};
 }

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<ComPtr<ID3D11ShaderResourceView>> CreateD3D11ShaderResourceView(
         uint64_t offset,
         uint64_t size) const;
     ResultOrError<ComPtr<ID3D11UnorderedAccessView1>> 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);
-
+    // Copy the buffer without checking if the buffer is initialized.
-    // The buffer object can be used as vertex, index, uniform, storage, or indirect buffer.
+    static MaybeError CopyInternal(CommandRecordingContext* commandContext,
-    ComPtr<ID3D11Buffer> mD3d11Buffer;
+                                   Buffer* source,
+                                   uint64_t sourceOffset,
+                                   size_t size,
+                                   Buffer* destination,
+                                   uint64_t destinationOffset);
+    // The buffer object for constant buffer usage.
+    ComPtr<ID3D11Buffer> mD3d11ConstantBuffer;
+    // The buffer object for non-constant buffer usages(e.g. storage buffer, vertex buffer, etc.)
+    ComPtr<ID3D11Buffer> mD3d11NonConstantBuffer;
+    bool mConstantBufferIsUpdated = true;
     uint8_t* mMappedData = nullptr;
 };
 

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<SetVertexBufferCmd>();
-                ID3D11Buffer* buffer = ToBackend(cmd->buffer)->GetD3D11Buffer();
+                ID3D11Buffer* buffer = ToBackend(cmd->buffer)->GetD3D11NonConstantBuffer();
                 vertexBufferTracker.OnSetVertexBuffer(cmd->slot, buffer, cmd->offset);
                 break;
             }

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,

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);
 

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",

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 <vector>
+
+#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 <typename T>
+    void CheckBuffer(ID3D11Buffer* buffer, std::vector<T> 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<ID3D11Buffer> 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<const T*>(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<uint8_t> 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<uint8_t> 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<uint32_t> data(kNumValues, 0x12345678);
+    uint64_t bufferSize = static_cast<uint64_t>(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<vec4u, 25>
+            }
+
+            @group(0) @binding(0) var<storage, read_write> 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<uint32_t> 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<vec4u, 25>
+            }
+
+            @group(0) @binding(0) var<uniform> src : Buf;
+            @group(0) @binding(1) var<storage, read_write> 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<uint32_t> 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());