From ebf183bde4e9c0b9f5a5a9836399e05ff5560033 Mon Sep 17 00:00:00 2001 From: Ken Rockot Date: Thu, 23 Sep 2021 00:15:19 +0000 Subject: [PATCH] Implement drawIndexedIndirect validation Every render pass which invokes DrawIndexedIndirect, either directly or through a RenderBundle execution, is now preceded immediately by at least one validation pass. All indirect buffer offests used with DII are validated, and their validated values are copied into a separate scratch buffer (or zeroed out there, in the case of validation failure). All encoded DII commands are rewritten to use the validated parameters instead of the original ones. Bug: dawn:809 Change-Id: I5eead937f19536f84f89e2c8e6fed7f18f0aee9f Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/63461 Commit-Queue: Ken Rockot Reviewed-by: Austin Eng --- src/common/NonCopyable.h | 3 + src/dawn_native/BUILD.gn | 7 + src/dawn_native/Buffer.cpp | 6 + src/dawn_native/CMakeLists.txt | 7 + src/dawn_native/CommandBuffer.cpp | 8 + src/dawn_native/CommandBuffer.h | 2 + src/dawn_native/CommandBufferStateTracker.cpp | 9 + src/dawn_native/CommandBufferStateTracker.h | 2 + src/dawn_native/CommandEncoder.cpp | 13 + src/dawn_native/CommandEncoder.h | 3 + src/dawn_native/Commands.cpp | 10 + src/dawn_native/Commands.h | 11 + src/dawn_native/ComputePassEncoder.cpp | 2 +- src/dawn_native/Device.cpp | 2 +- src/dawn_native/EncodingContext.cpp | 48 +- src/dawn_native/EncodingContext.h | 22 +- src/dawn_native/IndirectDrawMetadata.cpp | 193 +++++++ src/dawn_native/IndirectDrawMetadata.h | 112 ++++ .../IndirectDrawValidationEncoder.cpp | 397 ++++++++++++++ .../IndirectDrawValidationEncoder.h | 39 ++ src/dawn_native/InternalPipelineStore.cpp | 38 ++ src/dawn_native/InternalPipelineStore.h | 19 + src/dawn_native/RenderBundle.cpp | 8 +- src/dawn_native/RenderBundle.h | 6 +- src/dawn_native/RenderBundleEncoder.cpp | 3 +- src/dawn_native/RenderEncoderBase.cpp | 19 +- src/dawn_native/RenderEncoderBase.h | 2 + src/dawn_native/RenderPassEncoder.cpp | 8 +- src/dawn_native/ScratchBuffer.cpp | 47 ++ src/dawn_native/ScratchBuffer.h | 55 ++ src/dawn_native/d3d12/CommandBufferD3D12.cpp | 4 + src/dawn_native/metal/CommandBufferMTL.mm | 4 + src/dawn_native/opengl/CommandBufferGL.cpp | 4 + src/dawn_native/vulkan/CommandBufferVk.cpp | 4 + .../end2end/DrawIndexedIndirectTests.cpp | 507 +++++++++++++++++- .../validation/UnsafeAPIValidationTests.cpp | 81 --- 36 files changed, 1585 insertions(+), 120 deletions(-) create mode 100644 src/dawn_native/IndirectDrawMetadata.cpp create mode 100644 src/dawn_native/IndirectDrawMetadata.h create mode 100644 src/dawn_native/IndirectDrawValidationEncoder.cpp create mode 100644 src/dawn_native/IndirectDrawValidationEncoder.h create mode 100644 src/dawn_native/InternalPipelineStore.cpp create mode 100644 src/dawn_native/ScratchBuffer.cpp create mode 100644 src/dawn_native/ScratchBuffer.h diff --git a/src/common/NonCopyable.h b/src/common/NonCopyable.h index 61f15cabcf..2d217dfbad 100644 --- a/src/common/NonCopyable.h +++ b/src/common/NonCopyable.h @@ -21,6 +21,9 @@ class NonCopyable { constexpr NonCopyable() = default; ~NonCopyable() = default; + NonCopyable(NonCopyable&&) = default; + NonCopyable& operator=(NonCopyable&&) = default; + private: NonCopyable(const NonCopyable&) = delete; void operator=(const NonCopyable&) = delete; diff --git a/src/dawn_native/BUILD.gn b/src/dawn_native/BUILD.gn index b3b1f6e60b..6539d0bdf6 100644 --- a/src/dawn_native/BUILD.gn +++ b/src/dawn_native/BUILD.gn @@ -237,9 +237,14 @@ source_set("dawn_native_sources") { "Format.cpp", "Format.h", "Forward.h", + "IndirectDrawMetadata.cpp", + "IndirectDrawMetadata.h", + "IndirectDrawValidationEncoder.cpp", + "IndirectDrawValidationEncoder.h", "Instance.cpp", "Instance.h", "IntegerTypes.h", + "InternalPipelineStore.cpp", "InternalPipelineStore.h", "Limits.cpp", "Limits.h", @@ -286,6 +291,8 @@ source_set("dawn_native_sources") { "RingBufferAllocator.h", "Sampler.cpp", "Sampler.h", + "ScratchBuffer.cpp", + "ScratchBuffer.h", "ShaderModule.cpp", "ShaderModule.h", "StagingBuffer.cpp", diff --git a/src/dawn_native/Buffer.cpp b/src/dawn_native/Buffer.cpp index 7698c978aa..bd7064839d 100644 --- a/src/dawn_native/Buffer.cpp +++ b/src/dawn_native/Buffer.cpp @@ -147,6 +147,12 @@ namespace dawn_native { if (mUsage & wgpu::BufferUsage::QueryResolve) { mUsage |= kInternalStorageBuffer; } + + // We also add internal storage usage for Indirect buffers if validation is enabled, since + // validation involves binding them as storage buffers for use in a compute pass. + if ((mUsage & wgpu::BufferUsage::Indirect) && device->IsValidationEnabled()) { + mUsage |= kInternalStorageBuffer; + } } BufferBase::BufferBase(DeviceBase* device, diff --git a/src/dawn_native/CMakeLists.txt b/src/dawn_native/CMakeLists.txt index fb93f898fb..8e2b2ce9e5 100644 --- a/src/dawn_native/CMakeLists.txt +++ b/src/dawn_native/CMakeLists.txt @@ -98,6 +98,10 @@ target_sources(dawn_native PRIVATE "Extensions.h" "ExternalTexture.cpp" "ExternalTexture.h" + "IndirectDrawMetadata.cpp" + "IndirectDrawMetadata.h" + "IndirectDrawValidationEncoder.cpp" + "IndirectDrawValidationEncoder.h" "ObjectContentHasher.cpp" "ObjectContentHasher.h" "Format.cpp" @@ -105,6 +109,7 @@ target_sources(dawn_native PRIVATE "Forward.h" "Instance.cpp" "Instance.h" + "InternalPipelineStore.cpp" "InternalPipelineStore.h" "IntegerTypes.h" "Limits.cpp" @@ -150,6 +155,8 @@ target_sources(dawn_native PRIVATE "RingBufferAllocator.h" "Sampler.cpp" "Sampler.h" + "ScratchBuffer.cpp" + "ScratchBuffer.h" "ShaderModule.cpp" "ShaderModule.h" "StagingBuffer.cpp" diff --git a/src/dawn_native/CommandBuffer.cpp b/src/dawn_native/CommandBuffer.cpp index 43455c2178..1300cd1afc 100644 --- a/src/dawn_native/CommandBuffer.cpp +++ b/src/dawn_native/CommandBuffer.cpp @@ -38,6 +38,14 @@ namespace dawn_native { Destroy(); } + void CommandBufferBase::DoNextSetValidatedBufferLocationsInternal() { + SetValidatedBufferLocationsInternalCmd* cmd = + mCommands.NextCommand(); + for (const DeferredBufferLocationUpdate& update : cmd->updates) { + update.location->Set(update.buffer.Get(), update.offset); + } + } + // static CommandBufferBase* CommandBufferBase::MakeError(DeviceBase* device) { return new CommandBufferBase(device, ObjectBase::kError); diff --git a/src/dawn_native/CommandBuffer.h b/src/dawn_native/CommandBuffer.h index e90d320280..94159a10e9 100644 --- a/src/dawn_native/CommandBuffer.h +++ b/src/dawn_native/CommandBuffer.h @@ -44,6 +44,8 @@ namespace dawn_native { protected: ~CommandBufferBase(); + void DoNextSetValidatedBufferLocationsInternal(); + CommandIterator mCommands; private: diff --git a/src/dawn_native/CommandBufferStateTracker.cpp b/src/dawn_native/CommandBufferStateTracker.cpp index f3a6b47037..001be9105e 100644 --- a/src/dawn_native/CommandBufferStateTracker.cpp +++ b/src/dawn_native/CommandBufferStateTracker.cpp @@ -308,4 +308,13 @@ namespace dawn_native { PipelineLayoutBase* CommandBufferStateTracker::GetPipelineLayout() const { return mLastPipelineLayout; } + + wgpu::IndexFormat CommandBufferStateTracker::GetIndexFormat() const { + return mIndexFormat; + } + + uint64_t CommandBufferStateTracker::GetIndexBufferSize() const { + return mIndexBufferSize; + } + } // namespace dawn_native diff --git a/src/dawn_native/CommandBufferStateTracker.h b/src/dawn_native/CommandBufferStateTracker.h index 805a4fb144..0a6c587a98 100644 --- a/src/dawn_native/CommandBufferStateTracker.h +++ b/src/dawn_native/CommandBufferStateTracker.h @@ -47,6 +47,8 @@ namespace dawn_native { BindGroupBase* GetBindGroup(BindGroupIndex index) const; PipelineLayoutBase* GetPipelineLayout() const; + wgpu::IndexFormat GetIndexFormat() const; + uint64_t GetIndexBufferSize() const; private: MaybeError ValidateOperation(ValidationAspects requiredAspects); diff --git a/src/dawn_native/CommandEncoder.cpp b/src/dawn_native/CommandEncoder.cpp index a256382868..143ba7e744 100644 --- a/src/dawn_native/CommandEncoder.cpp +++ b/src/dawn_native/CommandEncoder.cpp @@ -508,6 +508,7 @@ namespace dawn_native { uint32_t width = 0; uint32_t height = 0; Ref attachmentState; + mEncodingContext.WillBeginRenderPass(); bool success = mEncodingContext.TryEncode(this, [&](CommandAllocator* allocator) -> MaybeError { uint32_t sampleCount = 0; @@ -922,6 +923,18 @@ namespace dawn_native { return commandBuffer.Detach(); } + void CommandEncoder::EncodeSetValidatedBufferLocationsInternal( + std::vector updates) { + ASSERT(GetDevice()->IsValidationEnabled()); + mEncodingContext.TryEncode(this, [&](CommandAllocator* allocator) -> MaybeError { + SetValidatedBufferLocationsInternalCmd* cmd = + allocator->Allocate( + Command::SetValidatedBufferLocationsInternal); + cmd->updates = std::move(updates); + return {}; + }); + } + ResultOrError> CommandEncoder::FinishInternal( const CommandBufferDescriptor* descriptor) { DeviceBase* device = GetDevice(); diff --git a/src/dawn_native/CommandEncoder.h b/src/dawn_native/CommandEncoder.h index 7eaca04dca..dbb33cfef5 100644 --- a/src/dawn_native/CommandEncoder.h +++ b/src/dawn_native/CommandEncoder.h @@ -76,6 +76,9 @@ namespace dawn_native { CommandBufferBase* APIFinish(const CommandBufferDescriptor* descriptor = nullptr); + void EncodeSetValidatedBufferLocationsInternal( + std::vector updates); + private: ResultOrError> FinishInternal( const CommandBufferDescriptor* descriptor); diff --git a/src/dawn_native/Commands.cpp b/src/dawn_native/Commands.cpp index c2bd0dde6d..ea122e7468 100644 --- a/src/dawn_native/Commands.cpp +++ b/src/dawn_native/Commands.cpp @@ -158,6 +158,12 @@ namespace dawn_native { cmd->~SetStencilReferenceCmd(); break; } + case Command::SetValidatedBufferLocationsInternal: { + SetValidatedBufferLocationsInternalCmd* cmd = + commands->NextCommand(); + cmd->~SetValidatedBufferLocationsInternalCmd(); + break; + } case Command::SetViewport: { SetViewportCmd* cmd = commands->NextCommand(); cmd->~SetViewportCmd(); @@ -313,6 +319,10 @@ namespace dawn_native { commands->NextCommand(); break; + case Command::SetValidatedBufferLocationsInternal: + commands->NextCommand(); + break; + case Command::SetViewport: commands->NextCommand(); break; diff --git a/src/dawn_native/Commands.h b/src/dawn_native/Commands.h index 333d19ba10..09acd10abf 100644 --- a/src/dawn_native/Commands.h +++ b/src/dawn_native/Commands.h @@ -63,6 +63,7 @@ namespace dawn_native { SetBlendConstant, SetBindGroup, SetIndexBuffer, + SetValidatedBufferLocationsInternal, SetVertexBuffer, WriteBuffer, WriteTimestamp, @@ -224,6 +225,16 @@ namespace dawn_native { uint32_t reference; }; + struct DeferredBufferLocationUpdate { + Ref location; + Ref buffer; + uint64_t offset; + }; + + struct SetValidatedBufferLocationsInternalCmd { + std::vector updates; + }; + struct SetViewportCmd { float x, y, width, height, minDepth, maxDepth; }; diff --git a/src/dawn_native/ComputePassEncoder.cpp b/src/dawn_native/ComputePassEncoder.cpp index dcc5df83c3..834fa6c629 100644 --- a/src/dawn_native/ComputePassEncoder.cpp +++ b/src/dawn_native/ComputePassEncoder.cpp @@ -67,7 +67,7 @@ namespace dawn_native { return {}; })) { - mEncodingContext->ExitPass(this, mUsageTracker.AcquireResourceUsage()); + mEncodingContext->ExitComputePass(this, mUsageTracker.AcquireResourceUsage()); } } diff --git a/src/dawn_native/Device.cpp b/src/dawn_native/Device.cpp index c209cd10b9..44287e0009 100644 --- a/src/dawn_native/Device.cpp +++ b/src/dawn_native/Device.cpp @@ -225,7 +225,7 @@ namespace dawn_native { mDynamicUploader = std::make_unique(this); mCallbackTaskManager = std::make_unique(); mDeprecationWarnings = std::make_unique(); - mInternalPipelineStore = std::make_unique(); + mInternalPipelineStore = std::make_unique(this); mPersistentCache = std::make_unique(this); ASSERT(GetPlatform() != nullptr); diff --git a/src/dawn_native/EncodingContext.cpp b/src/dawn_native/EncodingContext.cpp index 9e8812db83..9e7b960345 100644 --- a/src/dawn_native/EncodingContext.cpp +++ b/src/dawn_native/EncodingContext.cpp @@ -19,6 +19,7 @@ #include "dawn_native/Commands.h" #include "dawn_native/Device.h" #include "dawn_native/ErrorData.h" +#include "dawn_native/IndirectDrawValidationEncoder.h" #include "dawn_native/RenderBundleEncoder.h" namespace dawn_native { @@ -47,8 +48,9 @@ namespace dawn_native { } void EncodingContext::MoveToIterator() { + CommitCommands(std::move(mPendingCommands)); if (!mWasMovedToIterator) { - mIterator = CommandIterator(std::move(mAllocator)); + mIterator.AcquireCommandBlocks(std::move(mAllocators)); mWasMovedToIterator = true; } } @@ -67,6 +69,18 @@ namespace dawn_native { } } + void EncodingContext::WillBeginRenderPass() { + ASSERT(mCurrentEncoder == mTopLevelEncoder); + if (mDevice->IsValidationEnabled()) { + // When validation is enabled, we are going to want to capture all commands encoded + // between and including BeginRenderPassCmd and EndRenderPassCmd, and defer their + // sequencing util after we have a chance to insert any necessary validation + // commands. To support this we commit any current commands now, so that the + // impending BeginRenderPassCmd starts in a fresh CommandAllocator. + CommitCommands(std::move(mPendingCommands)); + } + } + void EncodingContext::EnterPass(const ObjectBase* passEncoder) { // Assert we're at the top level. ASSERT(mCurrentEncoder == mTopLevelEncoder); @@ -75,15 +89,34 @@ namespace dawn_native { mCurrentEncoder = passEncoder; } - void EncodingContext::ExitPass(const ObjectBase* passEncoder, RenderPassResourceUsage usages) { + MaybeError EncodingContext::ExitRenderPass(const ObjectBase* passEncoder, + RenderPassResourceUsageTracker usageTracker, + CommandEncoder* commandEncoder, + IndirectDrawMetadata indirectDrawMetadata) { ASSERT(mCurrentEncoder != mTopLevelEncoder); ASSERT(mCurrentEncoder == passEncoder); mCurrentEncoder = mTopLevelEncoder; - mRenderPassUsages.push_back(std::move(usages)); + + if (mDevice->IsValidationEnabled()) { + // With validation enabled, commands were committed just before BeginRenderPassCmd was + // encoded by our RenderPassEncoder (see WillBeginRenderPass above). This means + // mPendingCommands contains only the commands from BeginRenderPassCmd to + // EndRenderPassCmd, inclusive. Now we swap out this allocator with a fresh one to give + // the validation encoder a chance to insert its commands first. + CommandAllocator renderCommands = std::move(mPendingCommands); + DAWN_TRY(EncodeIndirectDrawValidationCommands(mDevice, commandEncoder, &usageTracker, + &indirectDrawMetadata)); + CommitCommands(std::move(mPendingCommands)); + CommitCommands(std::move(renderCommands)); + } + + mRenderPassUsages.push_back(usageTracker.AcquireResourceUsage()); + return {}; } - void EncodingContext::ExitPass(const ObjectBase* passEncoder, ComputePassResourceUsage usages) { + void EncodingContext::ExitComputePass(const ObjectBase* passEncoder, + ComputePassResourceUsage usages) { ASSERT(mCurrentEncoder != mTopLevelEncoder); ASSERT(mCurrentEncoder == passEncoder); @@ -126,6 +159,7 @@ namespace dawn_native { // if Finish() has been called. mCurrentEncoder = nullptr; mTopLevelEncoder = nullptr; + CommitCommands(std::move(mPendingCommands)); if (mError != nullptr) { return std::move(mError); @@ -136,6 +170,12 @@ namespace dawn_native { return {}; } + void EncodingContext::CommitCommands(CommandAllocator allocator) { + if (!allocator.IsEmpty()) { + mAllocators.push_back(std::move(allocator)); + } + } + bool EncodingContext::IsFinished() const { return mTopLevelEncoder == nullptr; } diff --git a/src/dawn_native/EncodingContext.h b/src/dawn_native/EncodingContext.h index b97e317abb..522d29ed63 100644 --- a/src/dawn_native/EncodingContext.h +++ b/src/dawn_native/EncodingContext.h @@ -18,6 +18,7 @@ #include "dawn_native/CommandAllocator.h" #include "dawn_native/Error.h" #include "dawn_native/ErrorData.h" +#include "dawn_native/IndirectDrawMetadata.h" #include "dawn_native/PassResourceUsageTracker.h" #include "dawn_native/dawn_platform.h" @@ -25,6 +26,7 @@ namespace dawn_native { + class CommandEncoder; class DeviceBase; class ObjectBase; @@ -69,13 +71,21 @@ namespace dawn_native { return false; } ASSERT(!mWasMovedToIterator); - return !ConsumedError(encodeFunction(&mAllocator)); + return !ConsumedError(encodeFunction(&mPendingCommands)); } + // Must be called prior to encoding a BeginRenderPassCmd. Note that it's OK to call this + // and then not actually call EnterPass+ExitRenderPass, for example if some other pass setup + // failed validation before the BeginRenderPassCmd could be encoded. + void WillBeginRenderPass(); + // Functions to set current encoder state void EnterPass(const ObjectBase* passEncoder); - void ExitPass(const ObjectBase* passEncoder, RenderPassResourceUsage usages); - void ExitPass(const ObjectBase* passEncoder, ComputePassResourceUsage usages); + MaybeError ExitRenderPass(const ObjectBase* passEncoder, + RenderPassResourceUsageTracker usageTracker, + CommandEncoder* commandEncoder, + IndirectDrawMetadata indirectDrawMetadata); + void ExitComputePass(const ObjectBase* passEncoder, ComputePassResourceUsage usages); MaybeError Finish(); const RenderPassUsages& GetRenderPassUsages() const; @@ -84,6 +94,8 @@ namespace dawn_native { ComputePassUsages AcquireComputePassUsages(); private: + void CommitCommands(CommandAllocator allocator); + bool IsFinished() const; void MoveToIterator(); @@ -104,7 +116,9 @@ namespace dawn_native { ComputePassUsages mComputePassUsages; bool mWereComputePassUsagesAcquired = false; - CommandAllocator mAllocator; + CommandAllocator mPendingCommands; + + std::vector mAllocators; CommandIterator mIterator; bool mWasMovedToIterator = false; bool mWereCommandsAcquired = false; diff --git a/src/dawn_native/IndirectDrawMetadata.cpp b/src/dawn_native/IndirectDrawMetadata.cpp new file mode 100644 index 0000000000..235935f08e --- /dev/null +++ b/src/dawn_native/IndirectDrawMetadata.cpp @@ -0,0 +1,193 @@ +// Copyright 2021 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 "dawn_native/IndirectDrawMetadata.h" + +#include "common/Constants.h" +#include "common/RefCounted.h" +#include "dawn_native/IndirectDrawValidationEncoder.h" +#include "dawn_native/RenderBundle.h" + +#include +#include + +namespace dawn_native { + + namespace { + + // In the unlikely scenario that indirect offsets used over a single buffer span more than + // this length of the buffer, we split the validation work into multiple batches. + constexpr uint64_t kMaxBatchOffsetRange = kMaxStorageBufferBindingSize - + kMinStorageBufferOffsetAlignment - + kDrawIndexedIndirectSize; + + } // namespace + + IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::IndexedIndirectBufferValidationInfo( + BufferBase* indirectBuffer) + : mIndirectBuffer(indirectBuffer) { + } + + void IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::AddIndexedIndirectDraw( + IndexedIndirectDraw draw) { + const uint64_t newOffset = draw.clientBufferOffset; + auto it = mBatches.begin(); + while (it != mBatches.end()) { + IndexedIndirectValidationBatch& batch = *it; + if (batch.draws.size() >= kMaxDrawCallsPerIndirectValidationBatch) { + // This batch is full. If its minOffset is to the right of the new offset, we can + // just insert a new batch here. + if (newOffset < batch.minOffset) { + break; + } + + // Otherwise keep looking. + ++it; + continue; + } + + if (newOffset >= batch.minOffset && newOffset <= batch.maxOffset) { + batch.draws.push_back(std::move(draw)); + return; + } + + if (newOffset < batch.minOffset && + batch.maxOffset - newOffset <= kMaxBatchOffsetRange) { + // We can extend this batch to the left in order to fit the new offset. + batch.minOffset = newOffset; + batch.draws.push_back(std::move(draw)); + return; + } + + if (newOffset > batch.maxOffset && + newOffset - batch.minOffset <= kMaxBatchOffsetRange) { + // We can extend this batch to the right in order to fit the new offset. + batch.maxOffset = newOffset; + batch.draws.push_back(std::move(draw)); + return; + } + + if (newOffset < batch.minOffset) { + // We want to insert a new batch just before this one. + break; + } + + ++it; + } + + IndexedIndirectValidationBatch newBatch; + newBatch.minOffset = newOffset; + newBatch.maxOffset = newOffset; + newBatch.draws.push_back(std::move(draw)); + + mBatches.insert(it, std::move(newBatch)); + } + + void IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::AddBatch( + const IndexedIndirectValidationBatch& newBatch) { + auto it = mBatches.begin(); + while (it != mBatches.end()) { + IndexedIndirectValidationBatch& batch = *it; + uint64_t min = std::min(newBatch.minOffset, batch.minOffset); + uint64_t max = std::max(newBatch.maxOffset, batch.maxOffset); + if (max - min <= kMaxBatchOffsetRange && batch.draws.size() + newBatch.draws.size() <= + kMaxDrawCallsPerIndirectValidationBatch) { + // This batch fits within the limits of an existing batch. Merge it. + batch.minOffset = min; + batch.maxOffset = max; + batch.draws.insert(batch.draws.end(), newBatch.draws.begin(), newBatch.draws.end()); + return; + } + + if (newBatch.minOffset < batch.minOffset) { + break; + } + + ++it; + } + mBatches.push_back(newBatch); + } + + const std::vector& + IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::GetBatches() const { + return mBatches; + } + + IndirectDrawMetadata::IndirectDrawMetadata() = default; + + IndirectDrawMetadata::~IndirectDrawMetadata() = default; + + IndirectDrawMetadata::IndirectDrawMetadata(IndirectDrawMetadata&&) = default; + + IndirectDrawMetadata& IndirectDrawMetadata::operator=(IndirectDrawMetadata&&) = default; + + IndirectDrawMetadata::IndexedIndirectBufferValidationInfoMap* + IndirectDrawMetadata::GetIndexedIndirectBufferValidationInfo() { + return &mIndexedIndirectBufferValidationInfo; + } + + void IndirectDrawMetadata::AddBundle(RenderBundleBase* bundle) { + auto result = mAddedBundles.insert(bundle); + if (!result.second) { + return; + } + + for (const auto& entry : + bundle->GetIndirectDrawMetadata().mIndexedIndirectBufferValidationInfo) { + const IndexedIndirectConfig& config = entry.first; + auto it = mIndexedIndirectBufferValidationInfo.lower_bound(config); + if (it != mIndexedIndirectBufferValidationInfo.end() && it->first == config) { + // We already have batches for the same config. Merge the new ones in. + for (const IndexedIndirectValidationBatch& batch : entry.second.GetBatches()) { + it->second.AddBatch(batch); + } + } else { + mIndexedIndirectBufferValidationInfo.emplace_hint(it, config, entry.second); + } + } + } + + void IndirectDrawMetadata::AddIndexedIndirectDraw( + wgpu::IndexFormat indexFormat, + uint64_t indexBufferSize, + BufferBase* indirectBuffer, + uint64_t indirectOffset, + BufferLocation* drawCmdIndirectBufferLocation) { + uint64_t numIndexBufferElements; + switch (indexFormat) { + case wgpu::IndexFormat::Uint16: + numIndexBufferElements = indexBufferSize / 2; + break; + case wgpu::IndexFormat::Uint32: + numIndexBufferElements = indexBufferSize / 4; + break; + case wgpu::IndexFormat::Undefined: + UNREACHABLE(); + } + + const IndexedIndirectConfig config(indirectBuffer, numIndexBufferElements); + auto it = mIndexedIndirectBufferValidationInfo.find(config); + if (it == mIndexedIndirectBufferValidationInfo.end()) { + auto result = mIndexedIndirectBufferValidationInfo.emplace( + config, IndexedIndirectBufferValidationInfo(indirectBuffer)); + it = result.first; + } + + IndexedIndirectDraw draw; + draw.clientBufferOffset = indirectOffset; + draw.bufferLocation = drawCmdIndirectBufferLocation; + it->second.AddIndexedIndirectDraw(std::move(draw)); + } + +} // namespace dawn_native diff --git a/src/dawn_native/IndirectDrawMetadata.h b/src/dawn_native/IndirectDrawMetadata.h new file mode 100644 index 0000000000..04c38e326b --- /dev/null +++ b/src/dawn_native/IndirectDrawMetadata.h @@ -0,0 +1,112 @@ +// Copyright 2021 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. + +#ifndef DAWNNATIVE_INDIRECTDRAWMETADATA_H_ +#define DAWNNATIVE_INDIRECTDRAWMETADATA_H_ + +#include "common/NonCopyable.h" +#include "common/RefCounted.h" +#include "dawn_native/Buffer.h" +#include "dawn_native/BufferLocation.h" +#include "dawn_native/CommandBufferStateTracker.h" +#include "dawn_native/Commands.h" + +#include +#include +#include +#include +#include + +namespace dawn_native { + + class RenderBundleBase; + + // Metadata corresponding to the validation requirements of a single render pass. This metadata + // is accumulated while its corresponding render pass is encoded, and is later used to encode + // validation commands to be inserted into the command buffer just before the render pass's own + // commands. + class IndirectDrawMetadata : public NonCopyable { + public: + struct IndexedIndirectDraw { + uint64_t clientBufferOffset; + Ref bufferLocation; + }; + + struct IndexedIndirectValidationBatch { + uint64_t minOffset; + uint64_t maxOffset; + std::vector draws; + }; + + // Tracks information about every draw call in this render pass which uses the same indirect + // buffer and the same-sized index buffer. Calls are grouped by indirect offset ranges so + // that validation work can be chunked efficiently if necessary. + class IndexedIndirectBufferValidationInfo { + public: + explicit IndexedIndirectBufferValidationInfo(BufferBase* indirectBuffer); + + // Logs a new drawIndexedIndirect call for the render pass. `cmd` is updated with an + // assigned (and deferred) buffer ref and relative offset before returning. + void AddIndexedIndirectDraw(IndexedIndirectDraw draw); + + // Adds draw calls from an already-computed batch, e.g. from a previously encoded + // RenderBundle. The added batch is merged into an existing batch if possible, otherwise + // it's added to mBatch. + void AddBatch(const IndexedIndirectValidationBatch& batch); + + const std::vector& GetBatches() const; + + private: + Ref mIndirectBuffer; + + // A list of information about validation batches that will need to be executed for the + // corresponding indirect buffer prior to a single render pass. These are kept sorted by + // minOffset and may overlap iff the number of offsets in one batch would otherwise + // exceed some large upper bound (roughly ~33M draw calls). + // + // Since the most common expected cases will overwhelmingly require only a single + // validation pass per render pass, this is optimized for efficient updates to a single + // batch rather than for efficient manipulation of a large number of batches. + std::vector mBatches; + }; + + // Combination of an indirect buffer reference, and the number of addressable index buffer + // elements at the time of a draw call. + using IndexedIndirectConfig = std::pair; + using IndexedIndirectBufferValidationInfoMap = + std::map; + + IndirectDrawMetadata(); + ~IndirectDrawMetadata(); + + IndirectDrawMetadata(IndirectDrawMetadata&&); + IndirectDrawMetadata& operator=(IndirectDrawMetadata&&); + + IndexedIndirectBufferValidationInfoMap* GetIndexedIndirectBufferValidationInfo(); + + void AddBundle(RenderBundleBase* bundle); + void AddIndexedIndirectDraw(wgpu::IndexFormat indexFormat, + uint64_t indexBufferSize, + BufferBase* indirectBuffer, + uint64_t indirectOffset, + BufferLocation* drawCmdIndirectBufferLocation); + + private: + IndexedIndirectBufferValidationInfoMap mIndexedIndirectBufferValidationInfo; + std::set mAddedBundles; + }; + +} // namespace dawn_native + +#endif // DAWNNATIVE_INDIRECTDRAWMETADATA_H_ diff --git a/src/dawn_native/IndirectDrawValidationEncoder.cpp b/src/dawn_native/IndirectDrawValidationEncoder.cpp new file mode 100644 index 0000000000..c3c2a043a9 --- /dev/null +++ b/src/dawn_native/IndirectDrawValidationEncoder.cpp @@ -0,0 +1,397 @@ +// Copyright 2021 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 "dawn_native/IndirectDrawValidationEncoder.h" + +#include "common/Constants.h" +#include "common/Math.h" +#include "dawn_native/BindGroup.h" +#include "dawn_native/BindGroupLayout.h" +#include "dawn_native/CommandEncoder.h" +#include "dawn_native/ComputePassEncoder.h" +#include "dawn_native/ComputePipeline.h" +#include "dawn_native/Device.h" +#include "dawn_native/InternalPipelineStore.h" +#include "dawn_native/Queue.h" + +#include +#include + +namespace dawn_native { + + namespace { + // NOTE: This must match the workgroup_size attribute on the compute entry point below. + constexpr uint64_t kWorkgroupSize = 64; + + // Equivalent to the BatchInfo struct defined in the shader below. + struct BatchInfo { + uint64_t numIndexBufferElements; + uint32_t numDraws; + uint32_t padding; + }; + + // TODO(https://crbug.com/dawn/1108): Propagate validation feedback from this shader in + // various failure modes. + static const char sRenderValidationShaderSource[] = R"( + let kNumIndirectParamsPerDrawCall = 5u; + + let kIndexCountEntry = 0u; + let kInstanceCountEntry = 1u; + let kFirstIndexEntry = 2u; + let kBaseVertexEntry = 3u; + let kFirstInstanceEntry = 4u; + + [[block]] struct BatchInfo { + numIndexBufferElementsLow: u32; + numIndexBufferElementsHigh: u32; + numDraws: u32; + padding: u32; + indirectOffsets: array; + }; + + [[block]] struct IndirectParams { + data: array; + }; + + [[group(0), binding(0)]] var batch: BatchInfo; + [[group(0), binding(1)]] var clientParams: IndirectParams; + [[group(0), binding(2)]] var validatedParams: IndirectParams; + + fn fail(drawIndex: u32) { + let index = drawIndex * kNumIndirectParamsPerDrawCall; + validatedParams.data[index + kIndexCountEntry] = 0u; + validatedParams.data[index + kInstanceCountEntry] = 0u; + validatedParams.data[index + kFirstIndexEntry] = 0u; + validatedParams.data[index + kBaseVertexEntry] = 0u; + validatedParams.data[index + kFirstInstanceEntry] = 0u; + } + + fn pass(drawIndex: u32) { + let vIndex = drawIndex * kNumIndirectParamsPerDrawCall; + let cIndex = batch.indirectOffsets[drawIndex]; + validatedParams.data[vIndex + kIndexCountEntry] = + clientParams.data[cIndex + kIndexCountEntry]; + validatedParams.data[vIndex + kInstanceCountEntry] = + clientParams.data[cIndex + kInstanceCountEntry]; + validatedParams.data[vIndex + kFirstIndexEntry] = + clientParams.data[cIndex + kFirstIndexEntry]; + validatedParams.data[vIndex + kBaseVertexEntry] = + clientParams.data[cIndex + kBaseVertexEntry]; + validatedParams.data[vIndex + kFirstInstanceEntry] = + clientParams.data[cIndex + kFirstInstanceEntry]; + } + + [[stage(compute), workgroup_size(64, 1, 1)]] + fn main([[builtin(global_invocation_id)]] id : vec3) { + if (id.x >= batch.numDraws) { + return; + } + + let clientIndex = batch.indirectOffsets[id.x]; + let firstInstance = clientParams.data[clientIndex + kFirstInstanceEntry]; + if (firstInstance != 0u) { + fail(id.x); + return; + } + + if (batch.numIndexBufferElementsHigh >= 2u) { + // firstIndex and indexCount are both u32. The maximum possible sum of these + // values is 0x1fffffffe, which is less than 0x200000000. Nothing to validate. + pass(id.x); + return; + } + + let firstIndex = clientParams.data[clientIndex + kFirstIndexEntry]; + if (batch.numIndexBufferElementsHigh == 0u && + batch.numIndexBufferElementsLow < firstIndex) { + fail(id.x); + return; + } + + // Note that this subtraction may underflow, but only when + // numIndexBufferElementsHigh is 1u. The result is still correct in that case. + let maxIndexCount = batch.numIndexBufferElementsLow - firstIndex; + let indexCount = clientParams.data[clientIndex + kIndexCountEntry]; + if (indexCount > maxIndexCount) { + fail(id.x); + return; + } + pass(id.x); + } + )"; + + ResultOrError GetOrCreateRenderValidationPipeline( + DeviceBase* device) { + InternalPipelineStore* store = device->GetInternalPipelineStore(); + + if (store->renderValidationPipeline == nullptr) { + // Create compute shader module if not cached before. + if (store->renderValidationShader == nullptr) { + ShaderModuleDescriptor descriptor; + ShaderModuleWGSLDescriptor wgslDesc; + wgslDesc.source = sRenderValidationShaderSource; + descriptor.nextInChain = reinterpret_cast(&wgslDesc); + DAWN_TRY_ASSIGN(store->renderValidationShader, + device->CreateShaderModule(&descriptor)); + } + + BindGroupLayoutEntry entries[3]; + entries[0].binding = 0; + entries[0].visibility = wgpu::ShaderStage::Compute; + entries[0].buffer.type = wgpu::BufferBindingType::ReadOnlyStorage; + entries[1].binding = 1; + entries[1].visibility = wgpu::ShaderStage::Compute; + entries[1].buffer.type = kInternalStorageBufferBinding; + entries[2].binding = 2; + entries[2].visibility = wgpu::ShaderStage::Compute; + entries[2].buffer.type = wgpu::BufferBindingType::Storage; + + BindGroupLayoutDescriptor bindGroupLayoutDescriptor; + bindGroupLayoutDescriptor.entryCount = 3; + bindGroupLayoutDescriptor.entries = entries; + Ref bindGroupLayout; + DAWN_TRY_ASSIGN(bindGroupLayout, + device->CreateBindGroupLayout(&bindGroupLayoutDescriptor, true)); + + PipelineLayoutDescriptor pipelineDescriptor; + pipelineDescriptor.bindGroupLayoutCount = 1; + pipelineDescriptor.bindGroupLayouts = &bindGroupLayout.Get(); + Ref pipelineLayout; + DAWN_TRY_ASSIGN(pipelineLayout, device->CreatePipelineLayout(&pipelineDescriptor)); + + ComputePipelineDescriptor computePipelineDescriptor = {}; + computePipelineDescriptor.layout = pipelineLayout.Get(); + computePipelineDescriptor.compute.module = store->renderValidationShader.Get(); + computePipelineDescriptor.compute.entryPoint = "main"; + + DAWN_TRY_ASSIGN(store->renderValidationPipeline, + device->CreateComputePipeline(&computePipelineDescriptor)); + } + + return store->renderValidationPipeline.Get(); + } + + size_t GetBatchDataSize(uint32_t numDraws) { + return sizeof(BatchInfo) + numDraws * sizeof(uint32_t); + } + + } // namespace + + const uint32_t kBatchDrawCallLimitByDispatchSize = + kMaxComputePerDimensionDispatchSize * kWorkgroupSize; + const uint32_t kBatchDrawCallLimitByStorageBindingSize = + (kMaxStorageBufferBindingSize - sizeof(BatchInfo)) / sizeof(uint32_t); + const uint32_t kMaxDrawCallsPerIndirectValidationBatch = + std::min(kBatchDrawCallLimitByDispatchSize, kBatchDrawCallLimitByStorageBindingSize); + + MaybeError EncodeIndirectDrawValidationCommands(DeviceBase* device, + CommandEncoder* commandEncoder, + RenderPassResourceUsageTracker* usageTracker, + IndirectDrawMetadata* indirectDrawMetadata) { + struct Batch { + const IndirectDrawMetadata::IndexedIndirectValidationBatch* metadata; + uint64_t numIndexBufferElements; + uint64_t dataBufferOffset; + uint64_t dataSize; + uint64_t clientIndirectOffset; + uint64_t clientIndirectSize; + uint64_t validatedParamsOffset; + uint64_t validatedParamsSize; + BatchInfo* batchInfo; + }; + + struct Pass { + BufferBase* clientIndirectBuffer; + uint64_t validatedParamsSize = 0; + uint64_t batchDataSize = 0; + std::unique_ptr batchData{nullptr, std::free}; + std::vector batches; + }; + + // First stage is grouping all batches into passes. We try to pack as many batches into a + // single pass as possible. Batches can be grouped together as long as they're validating + // data from the same indirect buffer, but they may still be split into multiple passes if + // the number of draw calls in a pass would exceed some (very high) upper bound. + uint64_t numTotalDrawCalls = 0; + size_t validatedParamsSize = 0; + std::vector passes; + IndirectDrawMetadata::IndexedIndirectBufferValidationInfoMap& bufferInfoMap = + *indirectDrawMetadata->GetIndexedIndirectBufferValidationInfo(); + if (bufferInfoMap.empty()) { + return {}; + } + + for (auto& entry : bufferInfoMap) { + const IndirectDrawMetadata::IndexedIndirectConfig& config = entry.first; + BufferBase* clientIndirectBuffer = config.first; + for (const IndirectDrawMetadata::IndexedIndirectValidationBatch& batch : + entry.second.GetBatches()) { + const uint64_t minOffsetFromAlignedBoundary = + batch.minOffset % kMinStorageBufferOffsetAlignment; + const uint64_t minOffsetAlignedDown = + batch.minOffset - minOffsetFromAlignedBoundary; + + Batch newBatch; + newBatch.metadata = &batch; + newBatch.numIndexBufferElements = config.second; + newBatch.dataSize = GetBatchDataSize(batch.draws.size()); + newBatch.clientIndirectOffset = minOffsetAlignedDown; + newBatch.clientIndirectSize = + batch.maxOffset + kDrawIndexedIndirectSize - minOffsetAlignedDown; + numTotalDrawCalls += batch.draws.size(); + + newBatch.validatedParamsSize = batch.draws.size() * kDrawIndexedIndirectSize; + newBatch.validatedParamsOffset = + Align(validatedParamsSize, kMinStorageBufferOffsetAlignment); + validatedParamsSize = newBatch.validatedParamsOffset + newBatch.validatedParamsSize; + if (validatedParamsSize > kMaxStorageBufferBindingSize) { + return DAWN_INTERNAL_ERROR("Too many drawIndexedIndirect calls to validate"); + } + + Pass* currentPass = passes.empty() ? nullptr : &passes.back(); + if (currentPass && currentPass->clientIndirectBuffer == clientIndirectBuffer) { + uint64_t nextBatchDataOffset = + Align(currentPass->batchDataSize, kMinStorageBufferOffsetAlignment); + uint64_t newPassBatchDataSize = nextBatchDataOffset + newBatch.dataSize; + if (newPassBatchDataSize <= kMaxStorageBufferBindingSize) { + // We can fit this batch in the current pass. + newBatch.dataBufferOffset = nextBatchDataOffset; + currentPass->batchDataSize = newPassBatchDataSize; + currentPass->batches.push_back(newBatch); + continue; + } + } + + // We need to start a new pass for this batch. + newBatch.dataBufferOffset = 0; + + Pass newPass; + newPass.clientIndirectBuffer = clientIndirectBuffer; + newPass.batchDataSize = newBatch.dataSize; + newPass.batches.push_back(newBatch); + passes.push_back(std::move(newPass)); + } + } + + auto* const store = device->GetInternalPipelineStore(); + ScratchBuffer& validatedParamsBuffer = store->scratchIndirectStorage; + ScratchBuffer& batchDataBuffer = store->scratchStorage; + + uint64_t requiredBatchDataBufferSize = 0; + for (const Pass& pass : passes) { + requiredBatchDataBufferSize = std::max(requiredBatchDataBufferSize, pass.batchDataSize); + } + DAWN_TRY(batchDataBuffer.EnsureCapacity(requiredBatchDataBufferSize)); + usageTracker->BufferUsedAs(batchDataBuffer.GetBuffer(), wgpu::BufferUsage::Storage); + + DAWN_TRY(validatedParamsBuffer.EnsureCapacity(validatedParamsSize)); + usageTracker->BufferUsedAs(validatedParamsBuffer.GetBuffer(), wgpu::BufferUsage::Indirect); + + // Now we allocate and populate host-side batch data to be copied to the GPU, and prepare to + // update all DrawIndexedIndirectCmd buffer references. + std::vector deferredBufferLocationUpdates; + deferredBufferLocationUpdates.reserve(numTotalDrawCalls); + for (Pass& pass : passes) { + // We use std::malloc here because it guarantees maximal scalar alignment. + pass.batchData = {std::malloc(pass.batchDataSize), std::free}; + memset(pass.batchData.get(), 0, pass.batchDataSize); + uint8_t* batchData = static_cast(pass.batchData.get()); + for (Batch& batch : pass.batches) { + batch.batchInfo = new (&batchData[batch.dataBufferOffset]) BatchInfo(); + batch.batchInfo->numIndexBufferElements = batch.numIndexBufferElements; + batch.batchInfo->numDraws = static_cast(batch.metadata->draws.size()); + + uint32_t* indirectOffsets = reinterpret_cast(batch.batchInfo + 1); + uint64_t validatedParamsOffset = batch.validatedParamsOffset; + for (const auto& draw : batch.metadata->draws) { + // The shader uses this to index an array of u32, hence the division by 4 bytes. + *indirectOffsets++ = static_cast( + (draw.clientBufferOffset - batch.clientIndirectOffset) / 4); + + DeferredBufferLocationUpdate deferredUpdate; + deferredUpdate.location = draw.bufferLocation; + deferredUpdate.buffer = validatedParamsBuffer.GetBuffer(); + deferredUpdate.offset = validatedParamsOffset; + deferredBufferLocationUpdates.push_back(std::move(deferredUpdate)); + + validatedParamsOffset += kDrawIndexedIndirectSize; + } + } + } + + ComputePipelineBase* pipeline; + DAWN_TRY_ASSIGN(pipeline, GetOrCreateRenderValidationPipeline(device)); + + Ref layout; + DAWN_TRY_ASSIGN(layout, pipeline->GetBindGroupLayout(0)); + + BindGroupEntry bindings[3]; + BindGroupEntry& bufferDataBinding = bindings[0]; + bufferDataBinding.binding = 0; + bufferDataBinding.buffer = batchDataBuffer.GetBuffer(); + + BindGroupEntry& clientIndirectBinding = bindings[1]; + clientIndirectBinding.binding = 1; + + BindGroupEntry& validatedParamsBinding = bindings[2]; + validatedParamsBinding.binding = 2; + validatedParamsBinding.buffer = validatedParamsBuffer.GetBuffer(); + + BindGroupDescriptor bindGroupDescriptor = {}; + bindGroupDescriptor.layout = layout.Get(); + bindGroupDescriptor.entryCount = 3; + bindGroupDescriptor.entries = bindings; + + // Finally, we can now encode our validation passes. Each pass first does a single + // WriteBuffer to get batch data over to the GPU, followed by a single compute pass. The + // compute pass encodes a separate SetBindGroup and Dispatch command for each batch. + commandEncoder->EncodeSetValidatedBufferLocationsInternal( + std::move(deferredBufferLocationUpdates)); + for (const Pass& pass : passes) { + commandEncoder->APIWriteBuffer(batchDataBuffer.GetBuffer(), 0, + static_cast(pass.batchData.get()), + pass.batchDataSize); + + // TODO(dawn:723): change to not use AcquireRef for reentrant object creation. + ComputePassDescriptor descriptor = {}; + Ref passEncoder = + AcquireRef(commandEncoder->APIBeginComputePass(&descriptor)); + passEncoder->APISetPipeline(pipeline); + + clientIndirectBinding.buffer = pass.clientIndirectBuffer; + + for (const Batch& batch : pass.batches) { + bufferDataBinding.offset = batch.dataBufferOffset; + bufferDataBinding.size = batch.dataSize; + clientIndirectBinding.offset = batch.clientIndirectOffset; + clientIndirectBinding.size = batch.clientIndirectSize; + validatedParamsBinding.offset = batch.validatedParamsOffset; + validatedParamsBinding.size = batch.validatedParamsSize; + + Ref bindGroup; + DAWN_TRY_ASSIGN(bindGroup, device->CreateBindGroup(&bindGroupDescriptor)); + + const uint32_t numDrawsRoundedUp = + (batch.batchInfo->numDraws + kWorkgroupSize - 1) / kWorkgroupSize; + passEncoder->APISetBindGroup(0, bindGroup.Get()); + passEncoder->APIDispatch(numDrawsRoundedUp); + } + + passEncoder->APIEndPass(); + } + + return {}; + } + +} // namespace dawn_native diff --git a/src/dawn_native/IndirectDrawValidationEncoder.h b/src/dawn_native/IndirectDrawValidationEncoder.h new file mode 100644 index 0000000000..bc62bf09b8 --- /dev/null +++ b/src/dawn_native/IndirectDrawValidationEncoder.h @@ -0,0 +1,39 @@ +// Copyright 2021 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. + +#ifndef DAWNNATIVE_INDIRECTDRAWVALIDATIONENCODER_H_ +#define DAWNNATIVE_INDIRECTDRAWVALIDATIONENCODER_H_ + +#include "dawn_native/Error.h" +#include "dawn_native/IndirectDrawMetadata.h" + +namespace dawn_native { + + class CommandEncoder; + class DeviceBase; + class RenderPassResourceUsageTracker; + + // The maximum number of draws call we can fit into a single validation batch. This is + // essentially limited by the number of indirect parameter blocks that can fit into the maximum + // allowed storage binding size (about 6.7M). + extern const uint32_t kMaxDrawCallsPerIndirectValidationBatch; + + MaybeError EncodeIndirectDrawValidationCommands(DeviceBase* device, + CommandEncoder* commandEncoder, + RenderPassResourceUsageTracker* usageTracker, + IndirectDrawMetadata* indirectDrawMetadata); + +} // namespace dawn_native + +#endif // DAWNNATIVE_INDIRECTDRAWVALIDATIONENCODER_H_ diff --git a/src/dawn_native/InternalPipelineStore.cpp b/src/dawn_native/InternalPipelineStore.cpp new file mode 100644 index 0000000000..edfd115f5e --- /dev/null +++ b/src/dawn_native/InternalPipelineStore.cpp @@ -0,0 +1,38 @@ +// Copyright 2021 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 "dawn_native/InternalPipelineStore.h" + +#include "dawn_native/ComputePipeline.h" +#include "dawn_native/Device.h" +#include "dawn_native/RenderPipeline.h" +#include "dawn_native/ShaderModule.h" + +#include + +namespace dawn_native { + + class RenderPipelineBase; + class ShaderModuleBase; + + InternalPipelineStore::InternalPipelineStore(DeviceBase* device) + : scratchStorage(device, wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Storage), + scratchIndirectStorage(device, + wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Indirect | + wgpu::BufferUsage::Storage) { + } + + InternalPipelineStore::~InternalPipelineStore() = default; + +} // namespace dawn_native diff --git a/src/dawn_native/InternalPipelineStore.h b/src/dawn_native/InternalPipelineStore.h index b3a7398f42..acf3b13dce 100644 --- a/src/dawn_native/InternalPipelineStore.h +++ b/src/dawn_native/InternalPipelineStore.h @@ -16,15 +16,23 @@ #define DAWNNATIVE_INTERNALPIPELINESTORE_H_ #include "dawn_native/ObjectBase.h" +#include "dawn_native/ScratchBuffer.h" #include "dawn_native/dawn_platform.h" #include namespace dawn_native { + + class DeviceBase; class RenderPipelineBase; class ShaderModuleBase; + // Every DeviceBase owns an InternalPipelineStore. This is a general-purpose cache for + // long-lived objects scoped to a device and used to support arbitrary pipeline operations. struct InternalPipelineStore { + explicit InternalPipelineStore(DeviceBase* device); + ~InternalPipelineStore(); + std::unordered_map> copyTextureForBrowserPipelines; @@ -34,7 +42,18 @@ namespace dawn_native { Ref timestampCS; Ref dummyFragmentShader; + + // A scratch buffer suitable for use as a copy destination and storage binding. + ScratchBuffer scratchStorage; + + // A scratch buffer suitable for use as a copy destination, storage binding, and indirect + // buffer for indirect dispatch or draw calls. + ScratchBuffer scratchIndirectStorage; + + Ref renderValidationPipeline; + Ref renderValidationShader; }; + } // namespace dawn_native #endif // DAWNNATIVE_INTERNALPIPELINESTORE_H_ diff --git a/src/dawn_native/RenderBundle.cpp b/src/dawn_native/RenderBundle.cpp index 028dde7e22..b17031b058 100644 --- a/src/dawn_native/RenderBundle.cpp +++ b/src/dawn_native/RenderBundle.cpp @@ -24,9 +24,11 @@ namespace dawn_native { RenderBundleBase::RenderBundleBase(RenderBundleEncoder* encoder, const RenderBundleDescriptor* descriptor, Ref attachmentState, - RenderPassResourceUsage resourceUsage) + RenderPassResourceUsage resourceUsage, + IndirectDrawMetadata indirectDrawMetadata) : ObjectBase(encoder->GetDevice(), kLabelNotImplemented), mCommands(encoder->AcquireCommands()), + mIndirectDrawMetadata(std::move(indirectDrawMetadata)), mAttachmentState(std::move(attachmentState)), mResourceUsage(std::move(resourceUsage)) { } @@ -58,4 +60,8 @@ namespace dawn_native { return mResourceUsage; } + const IndirectDrawMetadata& RenderBundleBase::GetIndirectDrawMetadata() { + return mIndirectDrawMetadata; + } + } // namespace dawn_native diff --git a/src/dawn_native/RenderBundle.h b/src/dawn_native/RenderBundle.h index f971ed6a36..40db9243f7 100644 --- a/src/dawn_native/RenderBundle.h +++ b/src/dawn_native/RenderBundle.h @@ -19,6 +19,7 @@ #include "dawn_native/AttachmentState.h" #include "dawn_native/CommandAllocator.h" #include "dawn_native/Error.h" +#include "dawn_native/IndirectDrawMetadata.h" #include "dawn_native/ObjectBase.h" #include "dawn_native/PassResourceUsage.h" @@ -36,7 +37,8 @@ namespace dawn_native { RenderBundleBase(RenderBundleEncoder* encoder, const RenderBundleDescriptor* descriptor, Ref attachmentState, - RenderPassResourceUsage resourceUsage); + RenderPassResourceUsage resourceUsage, + IndirectDrawMetadata indirectDrawMetadata); static RenderBundleBase* MakeError(DeviceBase* device); @@ -44,6 +46,7 @@ namespace dawn_native { const AttachmentState* GetAttachmentState() const; const RenderPassResourceUsage& GetResourceUsage() const; + const IndirectDrawMetadata& GetIndirectDrawMetadata(); protected: ~RenderBundleBase() override; @@ -52,6 +55,7 @@ namespace dawn_native { RenderBundleBase(DeviceBase* device, ErrorTag errorTag); CommandIterator mCommands; + IndirectDrawMetadata mIndirectDrawMetadata; Ref mAttachmentState; RenderPassResourceUsage mResourceUsage; }; diff --git a/src/dawn_native/RenderBundleEncoder.cpp b/src/dawn_native/RenderBundleEncoder.cpp index daff3eb33b..7ddda3154c 100644 --- a/src/dawn_native/RenderBundleEncoder.cpp +++ b/src/dawn_native/RenderBundleEncoder.cpp @@ -130,7 +130,8 @@ namespace dawn_native { DAWN_TRY(ValidateFinish(usages)); } - return new RenderBundleBase(this, descriptor, AcquireAttachmentState(), std::move(usages)); + return new RenderBundleBase(this, descriptor, AcquireAttachmentState(), std::move(usages), + std::move(mIndirectDrawMetadata)); } MaybeError RenderBundleEncoder::ValidateFinish(const RenderPassResourceUsage& usages) const { diff --git a/src/dawn_native/RenderEncoderBase.cpp b/src/dawn_native/RenderEncoderBase.cpp index 06b0f9b690..58849cfb2d 100644 --- a/src/dawn_native/RenderEncoderBase.cpp +++ b/src/dawn_native/RenderEncoderBase.cpp @@ -157,16 +157,6 @@ namespace dawn_native { DAWN_TRY(ValidateCanUseAs(indirectBuffer, wgpu::BufferUsage::Indirect)); DAWN_TRY(mCommandBufferState.ValidateCanDrawIndexed()); - // Indexed indirect draws need a compute-shader based validation check that the - // range of indices is contained inside the index buffer on Metal. Disallow them as - // unsafe until the validation is implemented. - if (GetDevice()->IsToggleEnabled(Toggle::DisallowUnsafeAPIs)) { - return DAWN_VALIDATION_ERROR( - "DrawIndexedIndirect is disallowed because it doesn't validate that the " - "index " - "range is valid yet."); - } - if (indirectOffset % 4 != 0) { return DAWN_VALIDATION_ERROR("Indirect offset must be a multiple of 4"); } @@ -179,7 +169,14 @@ namespace dawn_native { DrawIndexedIndirectCmd* cmd = allocator->Allocate(Command::DrawIndexedIndirect); - cmd->indirectBufferLocation = BufferLocation::New(indirectBuffer, indirectOffset); + if (IsValidationEnabled()) { + cmd->indirectBufferLocation = BufferLocation::New(); + mIndirectDrawMetadata.AddIndexedIndirectDraw( + mCommandBufferState.GetIndexFormat(), mCommandBufferState.GetIndexBufferSize(), + indirectBuffer, indirectOffset, cmd->indirectBufferLocation.Get()); + } else { + cmd->indirectBufferLocation = BufferLocation::New(indirectBuffer, indirectOffset); + } mUsageTracker.BufferUsedAs(indirectBuffer, wgpu::BufferUsage::Indirect); diff --git a/src/dawn_native/RenderEncoderBase.h b/src/dawn_native/RenderEncoderBase.h index 4976ee20cd..30b7a3ce88 100644 --- a/src/dawn_native/RenderEncoderBase.h +++ b/src/dawn_native/RenderEncoderBase.h @@ -18,6 +18,7 @@ #include "dawn_native/AttachmentState.h" #include "dawn_native/CommandBufferStateTracker.h" #include "dawn_native/Error.h" +#include "dawn_native/IndirectDrawMetadata.h" #include "dawn_native/PassResourceUsageTracker.h" #include "dawn_native/ProgrammablePassEncoder.h" @@ -64,6 +65,7 @@ namespace dawn_native { CommandBufferStateTracker mCommandBufferState; RenderPassResourceUsageTracker mUsageTracker; + IndirectDrawMetadata mIndirectDrawMetadata; private: Ref mAttachmentState; diff --git a/src/dawn_native/RenderPassEncoder.cpp b/src/dawn_native/RenderPassEncoder.cpp index b9cee4f7ad..250d064ca1 100644 --- a/src/dawn_native/RenderPassEncoder.cpp +++ b/src/dawn_native/RenderPassEncoder.cpp @@ -99,9 +99,11 @@ namespace dawn_native { } allocator->Allocate(Command::EndRenderPass); + DAWN_TRY(mEncodingContext->ExitRenderPass(this, std::move(mUsageTracker), + mCommandEncoder.Get(), + std::move(mIndirectDrawMetadata))); return {}; })) { - mEncodingContext->ExitPass(this, mUsageTracker.AcquireResourceUsage()); } } @@ -224,6 +226,10 @@ namespace dawn_native { mUsageTracker.AddRenderBundleTextureUsage(usages.textures[i], usages.textureUsages[i]); } + + if (IsValidationEnabled()) { + mIndirectDrawMetadata.AddBundle(renderBundles[i]); + } } return {}; diff --git a/src/dawn_native/ScratchBuffer.cpp b/src/dawn_native/ScratchBuffer.cpp new file mode 100644 index 0000000000..976214cb91 --- /dev/null +++ b/src/dawn_native/ScratchBuffer.cpp @@ -0,0 +1,47 @@ +// Copyright 2021 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 "dawn_native/ScratchBuffer.h" + +#include "dawn_native/Device.h" + +namespace dawn_native { + + ScratchBuffer::ScratchBuffer(DeviceBase* device, wgpu::BufferUsage usage) + : mDevice(device), mUsage(usage) { + } + + ScratchBuffer::~ScratchBuffer() = default; + + void ScratchBuffer::Reset() { + mBuffer = nullptr; + } + + MaybeError ScratchBuffer::EnsureCapacity(uint64_t capacity) { + if (!mBuffer.Get() || mBuffer->GetSize() < capacity) { + BufferDescriptor descriptor; + descriptor.size = capacity; + descriptor.usage = mUsage; + DAWN_TRY_ASSIGN(mBuffer, mDevice->CreateBuffer(&descriptor)); + mBuffer->SetIsDataInitialized(); + } + return {}; + } + + BufferBase* ScratchBuffer::GetBuffer() const { + ASSERT(mBuffer.Get() != nullptr); + return mBuffer.Get(); + } + +} // namespace dawn_native diff --git a/src/dawn_native/ScratchBuffer.h b/src/dawn_native/ScratchBuffer.h new file mode 100644 index 0000000000..7bb446dfa2 --- /dev/null +++ b/src/dawn_native/ScratchBuffer.h @@ -0,0 +1,55 @@ +// Copyright 2021 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. + +#ifndef DAWNNATIVE_SCRATCHBUFFER_H_ +#define DAWNNATIVE_SCRATCHBUFFER_H_ + +#include "common/RefCounted.h" +#include "dawn_native/Buffer.h" + +#include + +namespace dawn_native { + + class DeviceBase; + + // A ScratchBuffer is a lazily allocated and lazily grown GPU buffer for intermittent use by + // commands in the GPU queue. Note that scratch buffers are not zero-initialized, so users must + // be careful not to exposed uninitialized bytes to client shaders. + class ScratchBuffer { + public: + // Note that this object does not retain a reference to `device`, so `device` MUST outlive + // this object. + ScratchBuffer(DeviceBase* device, wgpu::BufferUsage usage); + ~ScratchBuffer(); + + // Resets this ScratchBuffer, guaranteeing that the next EnsureCapacity call allocates a + // fresh buffer. + void Reset(); + + // Ensures that this ScratchBuffer is backed by a buffer on `device` with at least + // `capacity` bytes of storage. + MaybeError EnsureCapacity(uint64_t capacity); + + BufferBase* GetBuffer() const; + + private: + DeviceBase* const mDevice; + const wgpu::BufferUsage mUsage; + Ref mBuffer; + }; + +} // namespace dawn_native + +#endif // DAWNNATIVE_SCRATCHBUFFER_H_ diff --git a/src/dawn_native/d3d12/CommandBufferD3D12.cpp b/src/dawn_native/d3d12/CommandBufferD3D12.cpp index 175e3dd411..4bdc3b0760 100644 --- a/src/dawn_native/d3d12/CommandBufferD3D12.cpp +++ b/src/dawn_native/d3d12/CommandBufferD3D12.cpp @@ -981,6 +981,10 @@ namespace dawn_native { namespace d3d12 { break; } + case Command::SetValidatedBufferLocationsInternal: + DoNextSetValidatedBufferLocationsInternal(); + break; + case Command::WriteBuffer: { WriteBufferCmd* write = mCommands.NextCommand(); const uint64_t offset = write->offset; diff --git a/src/dawn_native/metal/CommandBufferMTL.mm b/src/dawn_native/metal/CommandBufferMTL.mm index 4dd47a5704..eb9892b6af 100644 --- a/src/dawn_native/metal/CommandBufferMTL.mm +++ b/src/dawn_native/metal/CommandBufferMTL.mm @@ -987,6 +987,10 @@ namespace dawn_native { namespace metal { break; } + case Command::SetValidatedBufferLocationsInternal: + DoNextSetValidatedBufferLocationsInternal(); + break; + case Command::WriteBuffer: { WriteBufferCmd* write = mCommands.NextCommand(); const uint64_t offset = write->offset; diff --git a/src/dawn_native/opengl/CommandBufferGL.cpp b/src/dawn_native/opengl/CommandBufferGL.cpp index 099d590a68..f4bb61fa86 100644 --- a/src/dawn_native/opengl/CommandBufferGL.cpp +++ b/src/dawn_native/opengl/CommandBufferGL.cpp @@ -843,6 +843,10 @@ namespace dawn_native { namespace opengl { break; } + case Command::SetValidatedBufferLocationsInternal: + DoNextSetValidatedBufferLocationsInternal(); + break; + case Command::WriteBuffer: { WriteBufferCmd* write = mCommands.NextCommand(); uint64_t offset = write->offset; diff --git a/src/dawn_native/vulkan/CommandBufferVk.cpp b/src/dawn_native/vulkan/CommandBufferVk.cpp index cfc2a71fa8..36f6a39997 100644 --- a/src/dawn_native/vulkan/CommandBufferVk.cpp +++ b/src/dawn_native/vulkan/CommandBufferVk.cpp @@ -824,6 +824,10 @@ namespace dawn_native { namespace vulkan { break; } + case Command::SetValidatedBufferLocationsInternal: + DoNextSetValidatedBufferLocationsInternal(); + break; + case Command::WriteBuffer: { WriteBufferCmd* write = mCommands.NextCommand(); const uint64_t offset = write->offset; diff --git a/src/tests/end2end/DrawIndexedIndirectTests.cpp b/src/tests/end2end/DrawIndexedIndirectTests.cpp index 147d5122cb..0a8a83f941 100644 --- a/src/tests/end2end/DrawIndexedIndirectTests.cpp +++ b/src/tests/end2end/DrawIndexedIndirectTests.cpp @@ -14,6 +14,7 @@ #include "tests/DawnTest.h" +#include "utils/ComboRenderBundleEncoderDescriptor.h" #include "utils/ComboRenderPipelineDescriptor.h" #include "utils/WGPUHelpers.h" @@ -59,25 +60,26 @@ class DrawIndexedIndirectTest : public DawnTest { // Second quad: the first 3 vertices represent the top right triangle -1.0f, 1.0f, 0.0f, 1.0f, 1.0f, -1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, -1.0f, -1.0f, 0.0f, 1.0f}); - indexBuffer = utils::CreateBufferFromData( - device, wgpu::BufferUsage::Index, - {0, 1, 2, 0, 3, 1, - // The indices below are added to test negatve baseVertex - 0 + 4, 1 + 4, 2 + 4, 0 + 4, 3 + 4, 1 + 4}); } utils::BasicRenderPass renderPass; wgpu::RenderPipeline pipeline; wgpu::Buffer vertexBuffer; - wgpu::Buffer indexBuffer; - void Test(std::initializer_list bufferList, - uint64_t indexOffset, - uint64_t indirectOffset, - RGBA8 bottomLeftExpected, - RGBA8 topRightExpected) { - wgpu::Buffer indirectBuffer = - utils::CreateBufferFromData(device, wgpu::BufferUsage::Indirect, bufferList); + wgpu::Buffer CreateIndirectBuffer(std::initializer_list indirectParamList) { + return utils::CreateBufferFromData( + device, wgpu::BufferUsage::Indirect | wgpu::BufferUsage::Storage, indirectParamList); + } + + wgpu::Buffer CreateIndexBuffer(std::initializer_list indexList) { + return utils::CreateBufferFromData(device, wgpu::BufferUsage::Index, indexList); + } + + wgpu::CommandBuffer EncodeDrawCommands(std::initializer_list bufferList, + wgpu::Buffer indexBuffer, + uint64_t indexOffset, + uint64_t indirectOffset) { + wgpu::Buffer indirectBuffer = CreateIndirectBuffer(bufferList); wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); { @@ -89,12 +91,28 @@ class DrawIndexedIndirectTest : public DawnTest { pass.EndPass(); } - wgpu::CommandBuffer commands = encoder.Finish(); + return encoder.Finish(); + } + + void TestDraw(wgpu::CommandBuffer commands, RGBA8 bottomLeftExpected, RGBA8 topRightExpected) { queue.Submit(1, &commands); EXPECT_PIXEL_RGBA8_EQ(bottomLeftExpected, renderPass.color, 1, 3); EXPECT_PIXEL_RGBA8_EQ(topRightExpected, renderPass.color, 3, 1); } + + void Test(std::initializer_list bufferList, + uint64_t indexOffset, + uint64_t indirectOffset, + RGBA8 bottomLeftExpected, + RGBA8 topRightExpected) { + wgpu::Buffer indexBuffer = + CreateIndexBuffer({0, 1, 2, 0, 3, 1, + // The indices below are added to test negatve baseVertex + 0 + 4, 1 + 4, 2 + 4, 0 + 4, 3 + 4, 1 + 4}); + TestDraw(EncodeDrawCommands(bufferList, indexBuffer, indexOffset, indirectOffset), + bottomLeftExpected, topRightExpected); + } }; // The most basic DrawIndexed triangle draw. @@ -172,6 +190,467 @@ TEST_P(DrawIndexedIndirectTest, IndirectOffset) { Test({3, 1, 0, 4, 0, 3, 1, 3, 4, 0}, 0, 5 * sizeof(uint32_t), filled, notFilled); } +TEST_P(DrawIndexedIndirectTest, BasicValidation) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1}); + + // Test a draw with an excessive indexCount. Should draw nothing. + TestDraw(EncodeDrawCommands({7, 1, 0, 0, 0}, indexBuffer, 0, 0), notFilled, notFilled); + + // Test a draw with an excessive firstIndex. Should draw nothing. + TestDraw(EncodeDrawCommands({3, 1, 7, 0, 0}, indexBuffer, 0, 0), notFilled, notFilled); + + // Test a valid draw. Should draw only the second triangle. + TestDraw(EncodeDrawCommands({3, 1, 3, 0, 0}, indexBuffer, 0, 0), notFilled, filled); +} + +TEST_P(DrawIndexedIndirectTest, ValidateWithOffsets) { + // TODO(crbug.com/dawn/161): The GL/GLES backend doesn't support indirect index buffer offsets + // yet. + DAWN_SUPPRESS_TEST_IF(IsOpenGL() || IsOpenGLES()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + // Test that validation properly accounts for index buffer offset. + TestDraw(EncodeDrawCommands({3, 1, 0, 0, 0}, indexBuffer, 6 * sizeof(uint32_t), 0), filled, + notFilled); + TestDraw(EncodeDrawCommands({4, 1, 0, 0, 0}, indexBuffer, 6 * sizeof(uint32_t), 0), notFilled, + notFilled); + TestDraw(EncodeDrawCommands({3, 1, 4, 0, 0}, indexBuffer, 3 * sizeof(uint32_t), 0), notFilled, + notFilled); + + // Test that validation properly accounts for indirect buffer offset. + TestDraw( + EncodeDrawCommands({3, 1, 0, 0, 0, 1000, 1, 0, 0, 0}, indexBuffer, 0, 4 * sizeof(uint32_t)), + notFilled, notFilled); + TestDraw(EncodeDrawCommands({3, 1, 0, 0, 0, 1000, 1, 0, 0, 0}, indexBuffer, 0, 0), filled, + notFilled); +} + +TEST_P(DrawIndexedIndirectTest, ValidateMultiplePasses) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + // Test validation with multiple passes in a row. Namely this is exercising that scratch buffer + // data for use with a previous pass's validation commands is not overwritten before it can be + // used. + TestDraw(EncodeDrawCommands({10, 1, 0, 0, 0}, indexBuffer, 0, 0), notFilled, notFilled); + TestDraw(EncodeDrawCommands({6, 1, 0, 0, 0}, indexBuffer, 0, 0), filled, filled); + TestDraw(EncodeDrawCommands({4, 1, 6, 0, 0}, indexBuffer, 0, 0), notFilled, notFilled); + TestDraw(EncodeDrawCommands({3, 1, 6, 0, 0}, indexBuffer, 0, 0), filled, notFilled); + TestDraw(EncodeDrawCommands({3, 1, 3, 0, 0}, indexBuffer, 0, 0), notFilled, filled); + TestDraw(EncodeDrawCommands({6, 1, 3, 0, 0}, indexBuffer, 0, 0), filled, filled); + TestDraw(EncodeDrawCommands({6, 1, 6, 0, 0}, indexBuffer, 0, 0), notFilled, notFilled); +} + +TEST_P(DrawIndexedIndirectTest, ValidateMultipleDraws) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + // Validate multiple draw calls using the same index and indirect buffers as input, but with + // different indirect offsets. + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + { + wgpu::Buffer indirectBuffer = + CreateIndirectBuffer({3, 1, 3, 0, 0, 10, 1, 0, 0, 0, 3, 1, 6, 0, 0}); + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.SetPipeline(pipeline); + pass.SetVertexBuffer(0, vertexBuffer); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(indirectBuffer, 0); + pass.DrawIndexedIndirect(indirectBuffer, 20); + pass.DrawIndexedIndirect(indirectBuffer, 40); + pass.EndPass(); + } + + wgpu::CommandBuffer commands = encoder.Finish(); + + queue.Submit(1, &commands); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); + + // Validate multiple draw calls using the same indirect buffer but different index buffers as + // input. + encoder = device.CreateCommandEncoder(); + { + wgpu::Buffer indirectBuffer = + CreateIndirectBuffer({3, 1, 3, 0, 0, 10, 1, 0, 0, 0, 3, 1, 6, 0, 0}); + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.SetPipeline(pipeline); + pass.SetVertexBuffer(0, vertexBuffer); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(indirectBuffer, 0); + pass.SetIndexBuffer(CreateIndexBuffer({0, 3, 1, 0, 2, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(indirectBuffer, 20); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 2, 1}), + wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(indirectBuffer, 40); + pass.EndPass(); + } + commands = encoder.Finish(); + + queue.Submit(1, &commands); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); + + // Validate multiple draw calls using the same index buffer but different indirect buffers as + // input. + encoder = device.CreateCommandEncoder(); + { + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.SetPipeline(pipeline); + pass.SetVertexBuffer(0, vertexBuffer); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({3, 1, 3, 0, 0}), 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({10, 1, 0, 0, 0}), 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({3, 1, 6, 0, 0}), 0); + pass.EndPass(); + } + commands = encoder.Finish(); + + queue.Submit(1, &commands); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); + + // Validate multiple draw calls across different index and indirect buffers. + encoder = device.CreateCommandEncoder(); + { + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.SetPipeline(pipeline); + pass.SetVertexBuffer(0, vertexBuffer); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({3, 1, 3, 0, 0}), 0); + pass.SetIndexBuffer(CreateIndexBuffer({0, 1, 2, 0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({10, 1, 0, 0, 0}), 0); + pass.SetIndexBuffer(CreateIndexBuffer({0, 3, 1}), wgpu::IndexFormat::Uint32, 0); + pass.DrawIndexedIndirect(CreateIndirectBuffer({3, 1, 3, 0, 0}), 0); + pass.EndPass(); + } + commands = encoder.Finish(); + + queue.Submit(1, &commands); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); +} + +TEST_P(DrawIndexedIndirectTest, ValidateEncodeMultipleThenSubmitInOrder) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + wgpu::CommandBuffer commands[7]; + commands[0] = EncodeDrawCommands({10, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[1] = EncodeDrawCommands({6, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[2] = EncodeDrawCommands({4, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[3] = EncodeDrawCommands({3, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[4] = EncodeDrawCommands({3, 1, 3, 0, 0}, indexBuffer, 0, 0); + commands[5] = EncodeDrawCommands({6, 1, 3, 0, 0}, indexBuffer, 0, 0); + commands[6] = EncodeDrawCommands({6, 1, 6, 0, 0}, indexBuffer, 0, 0); + + TestDraw(commands[0], notFilled, notFilled); + TestDraw(commands[1], filled, filled); + TestDraw(commands[2], notFilled, notFilled); + TestDraw(commands[3], filled, notFilled); + TestDraw(commands[4], notFilled, filled); + TestDraw(commands[5], filled, filled); + TestDraw(commands[6], notFilled, notFilled); +} + +TEST_P(DrawIndexedIndirectTest, ValidateEncodeMultipleThenSubmitAtOnce) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // TODO(crbug.com/dawn/1124): Fails on Intel+Vulkan+Windows for drivers + // older than 27.20.100.8587, which bots are actively using. + DAWN_SUPPRESS_TEST_IF(IsIntel() && IsVulkan() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + wgpu::CommandBuffer commands[5]; + commands[0] = EncodeDrawCommands({10, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[1] = EncodeDrawCommands({6, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[2] = EncodeDrawCommands({4, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[3] = EncodeDrawCommands({3, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[4] = EncodeDrawCommands({3, 1, 3, 0, 0}, indexBuffer, 0, 0); + + queue.Submit(5, commands); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); +} + +TEST_P(DrawIndexedIndirectTest, ValidateEncodeMultipleThenSubmitOutOfOrder) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + wgpu::CommandBuffer commands[7]; + commands[0] = EncodeDrawCommands({10, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[1] = EncodeDrawCommands({6, 1, 0, 0, 0}, indexBuffer, 0, 0); + commands[2] = EncodeDrawCommands({4, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[3] = EncodeDrawCommands({3, 1, 6, 0, 0}, indexBuffer, 0, 0); + commands[4] = EncodeDrawCommands({3, 1, 3, 0, 0}, indexBuffer, 0, 0); + commands[5] = EncodeDrawCommands({6, 1, 3, 0, 0}, indexBuffer, 0, 0); + commands[6] = EncodeDrawCommands({6, 1, 6, 0, 0}, indexBuffer, 0, 0); + + TestDraw(commands[6], notFilled, notFilled); + TestDraw(commands[5], filled, filled); + TestDraw(commands[4], notFilled, filled); + TestDraw(commands[3], filled, notFilled); + TestDraw(commands[2], notFilled, notFilled); + TestDraw(commands[1], filled, filled); + TestDraw(commands[0], notFilled, notFilled); +} + +TEST_P(DrawIndexedIndirectTest, ValidateWithBundlesInSamePass) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indirectBuffer = + CreateIndirectBuffer({3, 1, 3, 0, 0, 10, 1, 0, 0, 0, 3, 1, 6, 0, 0}); + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + std::vector bundles; + { + utils::ComboRenderBundleEncoderDescriptor desc = {}; + desc.colorFormatsCount = 1; + desc.cColorFormats[0] = wgpu::TextureFormat::RGBA8Unorm; + wgpu::RenderBundleEncoder bundleEncoder = device.CreateRenderBundleEncoder(&desc); + bundleEncoder.SetPipeline(pipeline); + bundleEncoder.SetVertexBuffer(0, vertexBuffer); + bundleEncoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); + bundleEncoder.DrawIndexedIndirect(indirectBuffer, 20); + bundles.push_back(bundleEncoder.Finish()); + } + { + utils::ComboRenderBundleEncoderDescriptor desc = {}; + desc.colorFormatsCount = 1; + desc.cColorFormats[0] = wgpu::TextureFormat::RGBA8Unorm; + wgpu::RenderBundleEncoder bundleEncoder = device.CreateRenderBundleEncoder(&desc); + bundleEncoder.SetPipeline(pipeline); + bundleEncoder.SetVertexBuffer(0, vertexBuffer); + bundleEncoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); + bundleEncoder.DrawIndexedIndirect(indirectBuffer, 40); + bundles.push_back(bundleEncoder.Finish()); + } + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + { + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.ExecuteBundles(bundles.size(), bundles.data()); + pass.EndPass(); + } + wgpu::CommandBuffer commands = encoder.Finish(); + + queue.Submit(1, &commands); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 3, 1); +} + +TEST_P(DrawIndexedIndirectTest, ValidateWithBundlesInDifferentPasses) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows only. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indirectBuffer = + CreateIndirectBuffer({3, 1, 3, 0, 0, 10, 1, 0, 0, 0, 3, 1, 6, 0, 0}); + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1, 0, 1, 2}); + + wgpu::CommandBuffer commands[2]; + { + wgpu::RenderBundle bundle; + utils::ComboRenderBundleEncoderDescriptor desc = {}; + desc.colorFormatsCount = 1; + desc.cColorFormats[0] = wgpu::TextureFormat::RGBA8Unorm; + wgpu::RenderBundleEncoder bundleEncoder = device.CreateRenderBundleEncoder(&desc); + bundleEncoder.SetPipeline(pipeline); + bundleEncoder.SetVertexBuffer(0, vertexBuffer); + bundleEncoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); + bundleEncoder.DrawIndexedIndirect(indirectBuffer, 20); + bundle = bundleEncoder.Finish(); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + renderPass.renderPassInfo.cColorAttachments[0].loadOp = wgpu::LoadOp::Load; + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.ExecuteBundles(1, &bundle); + pass.EndPass(); + + commands[0] = encoder.Finish(); + } + + { + wgpu::RenderBundle bundle; + utils::ComboRenderBundleEncoderDescriptor desc = {}; + desc.colorFormatsCount = 1; + desc.cColorFormats[0] = wgpu::TextureFormat::RGBA8Unorm; + wgpu::RenderBundleEncoder bundleEncoder = device.CreateRenderBundleEncoder(&desc); + bundleEncoder.SetPipeline(pipeline); + bundleEncoder.SetVertexBuffer(0, vertexBuffer); + bundleEncoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); + bundleEncoder.DrawIndexedIndirect(indirectBuffer, 40); + bundle = bundleEncoder.Finish(); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + renderPass.renderPassInfo.cColorAttachments[0].loadOp = wgpu::LoadOp::Clear; + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.ExecuteBundles(1, &bundle); + pass.EndPass(); + + commands[1] = encoder.Finish(); + } + + queue.Submit(1, &commands[1]); + queue.Submit(1, &commands[0]); + + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(notFilled, renderPass.color, 3, 1); +} + +TEST_P(DrawIndexedIndirectTest, ValidateReusedBundleWithChangingParams) { + // TODO(crbug.com/dawn/789): Test is failing under SwANGLE on Windows. + DAWN_SUPPRESS_TEST_IF(IsANGLE() && IsWindows()); + + // TODO(crbug.com/dawn/1124): Fails on Intel+Vulkan+Windows for drivers + // older than 27.20.100.8587, which bots are actively using. + DAWN_SUPPRESS_TEST_IF(IsIntel() && IsVulkan() && IsWindows()); + + // It doesn't make sense to test invalid inputs when validation is disabled. + DAWN_SUPPRESS_TEST_IF(HasToggleEnabled("skip_validation")); + + RGBA8 filled(0, 255, 0, 255); + // RGBA8 notFilled(0, 0, 0, 0); + + wgpu::Buffer indirectBuffer = CreateIndirectBuffer({0, 0, 0, 0, 0}); + wgpu::Buffer indexBuffer = CreateIndexBuffer({0, 1, 2, 0, 3, 1}); + + // Encode a single bundle that always uses indirectBuffer offset 0 for its params. + wgpu::RenderBundle bundle; + utils::ComboRenderBundleEncoderDescriptor desc = {}; + desc.colorFormatsCount = 1; + desc.cColorFormats[0] = wgpu::TextureFormat::RGBA8Unorm; + wgpu::RenderBundleEncoder bundleEncoder = device.CreateRenderBundleEncoder(&desc); + bundleEncoder.SetPipeline(pipeline); + bundleEncoder.SetVertexBuffer(0, vertexBuffer); + bundleEncoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32, 0); + bundleEncoder.DrawIndexedIndirect(indirectBuffer, 0); + bundle = bundleEncoder.Finish(); + + wgpu::ShaderModule paramWriterModule = utils::CreateShaderModule(device, + R"( + [[block]] struct Input { firstIndex: u32; }; + [[block]] struct Params { + indexCount: u32; + instanceCount: u32; + firstIndex: u32; + }; + [[group(0), binding(0)]] var input: Input; + [[group(0), binding(1)]] var params: Params; + [[stage(compute), workgroup_size(1)]] fn main() { + params.indexCount = 3u; + params.instanceCount = 1u; + params.firstIndex = input.firstIndex; + } + )"); + + wgpu::ComputePipelineDescriptor computeDesc; + computeDesc.compute.module = paramWriterModule; + computeDesc.compute.entryPoint = "main"; + wgpu::ComputePipeline computePipeline = device.CreateComputePipeline(&computeDesc); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + + auto encodeComputePassToUpdateFirstIndex = [&](uint32_t newFirstIndex) { + wgpu::Buffer input = utils::CreateBufferFromData( + device, wgpu::BufferUsage::Uniform, {newFirstIndex}); + wgpu::BindGroup bindGroup = utils::MakeBindGroup( + device, computePipeline.GetBindGroupLayout(0), + {{0, input, 0, sizeof(uint32_t)}, {1, indirectBuffer, 0, 5 * sizeof(uint32_t)}}); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(computePipeline); + pass.SetBindGroup(0, bindGroup); + pass.Dispatch(1); + pass.EndPass(); + }; + + auto encodeRenderPassToExecuteBundle = [&](wgpu::LoadOp colorLoadOp) { + renderPass.renderPassInfo.cColorAttachments[0].loadOp = colorLoadOp; + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.ExecuteBundles(1, &bundle); + pass.EndPass(); + }; + + encodeComputePassToUpdateFirstIndex(0); + encodeRenderPassToExecuteBundle(wgpu::LoadOp::Clear); + encodeComputePassToUpdateFirstIndex(3); + encodeRenderPassToExecuteBundle(wgpu::LoadOp::Load); + encodeComputePassToUpdateFirstIndex(6); + encodeRenderPassToExecuteBundle(wgpu::LoadOp::Load); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 1, 3); + EXPECT_PIXEL_RGBA8_EQ(filled, renderPass.color, 3, 1); +} + DAWN_INSTANTIATE_TEST(DrawIndexedIndirectTest, D3D12Backend(), MetalBackend(), diff --git a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp index abac251e48..ffdb128240 100644 --- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp +++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp @@ -28,87 +28,6 @@ class UnsafeAPIValidationTest : public ValidationTest { } }; -// Check that DrawIndexedIndirect is disallowed as part of unsafe APIs. -TEST_F(UnsafeAPIValidationTest, DrawIndexedIndirectDisallowed) { - // Create the index and indirect buffers. - wgpu::BufferDescriptor indexBufferDesc; - indexBufferDesc.size = 4; - indexBufferDesc.usage = wgpu::BufferUsage::Index; - wgpu::Buffer indexBuffer = device.CreateBuffer(&indexBufferDesc); - - wgpu::BufferDescriptor indirectBufferDesc; - indirectBufferDesc.size = 64; - indirectBufferDesc.usage = wgpu::BufferUsage::Indirect; - wgpu::Buffer indirectBuffer = device.CreateBuffer(&indirectBufferDesc); - - // The RenderPassDescriptor, RenderBundleDescriptor and pipeline for all sub-tests below. - DummyRenderPass renderPass(device); - - utils::ComboRenderBundleEncoderDescriptor bundleDesc = {}; - bundleDesc.colorFormatsCount = 1; - bundleDesc.cColorFormats[0] = renderPass.attachmentFormat; - - utils::ComboRenderPipelineDescriptor desc; - desc.vertex.module = utils::CreateShaderModule( - device, - R"([[stage(vertex)]] fn main() -> [[builtin(position)]] vec4 { - return vec4(); - })"); - desc.cFragment.module = utils::CreateShaderModule(device, "[[stage(fragment)]] fn main() {}"); - desc.cTargets[0].writeMask = wgpu::ColorWriteMask::None; - wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&desc); - - // Control cases: DrawIndirect and DrawIndexed are allowed inside a render pass. - { - wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); - wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass); - pass.SetPipeline(pipeline); - - pass.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32); - pass.DrawIndexed(1); - - pass.DrawIndirect(indirectBuffer, 0); - pass.EndPass(); - encoder.Finish(); - } - - // Control case: DrawIndirect and DrawIndexed are allowed inside a render bundle. - { - wgpu::RenderBundleEncoder encoder = device.CreateRenderBundleEncoder(&bundleDesc); - encoder.SetPipeline(pipeline); - - encoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32); - encoder.DrawIndexed(1); - - encoder.DrawIndirect(indirectBuffer, 0); - encoder.Finish(); - } - - // Error case, DrawIndexedIndirect is disallowed inside a render pass. - { - wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); - wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass); - - pass.SetPipeline(pipeline); - pass.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32); - pass.DrawIndexedIndirect(indirectBuffer, 0); - - pass.EndPass(); - ASSERT_DEVICE_ERROR(encoder.Finish()); - } - - // Error case, DrawIndexedIndirect is disallowed inside a render bundle. - { - wgpu::RenderBundleEncoder encoder = device.CreateRenderBundleEncoder(&bundleDesc); - - encoder.SetPipeline(pipeline); - encoder.SetIndexBuffer(indexBuffer, wgpu::IndexFormat::Uint32); - encoder.DrawIndexedIndirect(indirectBuffer, 0); - - ASSERT_DEVICE_ERROR(encoder.Finish()); - } -} - // Check that DispatchIndirect is disallowed as part of unsafe APIs. TEST_F(UnsafeAPIValidationTest, DispatchIndirectDisallowed) { // Create the index and indirect buffers.