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 <rockot@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
diff --git a/src/common/NonCopyable.h b/src/common/NonCopyable.h
index 61f15ca..2d217df 100644
--- a/src/common/NonCopyable.h
+++ b/src/common/NonCopyable.h
@@ -21,6 +21,9 @@
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 b3b1f6e..6539d0b 100644
--- a/src/dawn_native/BUILD.gn
+++ b/src/dawn_native/BUILD.gn
@@ -237,9 +237,14 @@
"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 @@
"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 7698c97..bd70648 100644
--- a/src/dawn_native/Buffer.cpp
+++ b/src/dawn_native/Buffer.cpp
@@ -147,6 +147,12 @@
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 fb93f89..8e2b2ce 100644
--- a/src/dawn_native/CMakeLists.txt
+++ b/src/dawn_native/CMakeLists.txt
@@ -98,6 +98,10 @@
"Extensions.h"
"ExternalTexture.cpp"
"ExternalTexture.h"
+ "IndirectDrawMetadata.cpp"
+ "IndirectDrawMetadata.h"
+ "IndirectDrawValidationEncoder.cpp"
+ "IndirectDrawValidationEncoder.h"
"ObjectContentHasher.cpp"
"ObjectContentHasher.h"
"Format.cpp"
@@ -105,6 +109,7 @@
"Forward.h"
"Instance.cpp"
"Instance.h"
+ "InternalPipelineStore.cpp"
"InternalPipelineStore.h"
"IntegerTypes.h"
"Limits.cpp"
@@ -150,6 +155,8 @@
"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 43455c2..1300cd1 100644
--- a/src/dawn_native/CommandBuffer.cpp
+++ b/src/dawn_native/CommandBuffer.cpp
@@ -38,6 +38,14 @@
Destroy();
}
+ void CommandBufferBase::DoNextSetValidatedBufferLocationsInternal() {
+ SetValidatedBufferLocationsInternalCmd* cmd =
+ mCommands.NextCommand<SetValidatedBufferLocationsInternalCmd>();
+ 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 e90d320..94159a1 100644
--- a/src/dawn_native/CommandBuffer.h
+++ b/src/dawn_native/CommandBuffer.h
@@ -44,6 +44,8 @@
protected:
~CommandBufferBase();
+ void DoNextSetValidatedBufferLocationsInternal();
+
CommandIterator mCommands;
private:
diff --git a/src/dawn_native/CommandBufferStateTracker.cpp b/src/dawn_native/CommandBufferStateTracker.cpp
index f3a6b47..001be91 100644
--- a/src/dawn_native/CommandBufferStateTracker.cpp
+++ b/src/dawn_native/CommandBufferStateTracker.cpp
@@ -308,4 +308,13 @@
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 805a4fb..0a6c587 100644
--- a/src/dawn_native/CommandBufferStateTracker.h
+++ b/src/dawn_native/CommandBufferStateTracker.h
@@ -47,6 +47,8 @@
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 a256382..143ba7e 100644
--- a/src/dawn_native/CommandEncoder.cpp
+++ b/src/dawn_native/CommandEncoder.cpp
@@ -508,6 +508,7 @@
uint32_t width = 0;
uint32_t height = 0;
Ref<AttachmentState> attachmentState;
+ mEncodingContext.WillBeginRenderPass();
bool success =
mEncodingContext.TryEncode(this, [&](CommandAllocator* allocator) -> MaybeError {
uint32_t sampleCount = 0;
@@ -922,6 +923,18 @@
return commandBuffer.Detach();
}
+ void CommandEncoder::EncodeSetValidatedBufferLocationsInternal(
+ std::vector<DeferredBufferLocationUpdate> updates) {
+ ASSERT(GetDevice()->IsValidationEnabled());
+ mEncodingContext.TryEncode(this, [&](CommandAllocator* allocator) -> MaybeError {
+ SetValidatedBufferLocationsInternalCmd* cmd =
+ allocator->Allocate<SetValidatedBufferLocationsInternalCmd>(
+ Command::SetValidatedBufferLocationsInternal);
+ cmd->updates = std::move(updates);
+ return {};
+ });
+ }
+
ResultOrError<Ref<CommandBufferBase>> CommandEncoder::FinishInternal(
const CommandBufferDescriptor* descriptor) {
DeviceBase* device = GetDevice();
diff --git a/src/dawn_native/CommandEncoder.h b/src/dawn_native/CommandEncoder.h
index 7eaca04..dbb33cf 100644
--- a/src/dawn_native/CommandEncoder.h
+++ b/src/dawn_native/CommandEncoder.h
@@ -76,6 +76,9 @@
CommandBufferBase* APIFinish(const CommandBufferDescriptor* descriptor = nullptr);
+ void EncodeSetValidatedBufferLocationsInternal(
+ std::vector<DeferredBufferLocationUpdate> updates);
+
private:
ResultOrError<Ref<CommandBufferBase>> FinishInternal(
const CommandBufferDescriptor* descriptor);
diff --git a/src/dawn_native/Commands.cpp b/src/dawn_native/Commands.cpp
index c2bd0dd..ea122e7 100644
--- a/src/dawn_native/Commands.cpp
+++ b/src/dawn_native/Commands.cpp
@@ -158,6 +158,12 @@
cmd->~SetStencilReferenceCmd();
break;
}
+ case Command::SetValidatedBufferLocationsInternal: {
+ SetValidatedBufferLocationsInternalCmd* cmd =
+ commands->NextCommand<SetValidatedBufferLocationsInternalCmd>();
+ cmd->~SetValidatedBufferLocationsInternalCmd();
+ break;
+ }
case Command::SetViewport: {
SetViewportCmd* cmd = commands->NextCommand<SetViewportCmd>();
cmd->~SetViewportCmd();
@@ -313,6 +319,10 @@
commands->NextCommand<SetStencilReferenceCmd>();
break;
+ case Command::SetValidatedBufferLocationsInternal:
+ commands->NextCommand<SetValidatedBufferLocationsInternalCmd>();
+ break;
+
case Command::SetViewport:
commands->NextCommand<SetViewportCmd>();
break;
diff --git a/src/dawn_native/Commands.h b/src/dawn_native/Commands.h
index 333d19b..09acd10 100644
--- a/src/dawn_native/Commands.h
+++ b/src/dawn_native/Commands.h
@@ -63,6 +63,7 @@
SetBlendConstant,
SetBindGroup,
SetIndexBuffer,
+ SetValidatedBufferLocationsInternal,
SetVertexBuffer,
WriteBuffer,
WriteTimestamp,
@@ -224,6 +225,16 @@
uint32_t reference;
};
+ struct DeferredBufferLocationUpdate {
+ Ref<BufferLocation> location;
+ Ref<BufferBase> buffer;
+ uint64_t offset;
+ };
+
+ struct SetValidatedBufferLocationsInternalCmd {
+ std::vector<DeferredBufferLocationUpdate> 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 dcc5df8..834fa6c 100644
--- a/src/dawn_native/ComputePassEncoder.cpp
+++ b/src/dawn_native/ComputePassEncoder.cpp
@@ -67,7 +67,7 @@
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 c209cd1..44287e0 100644
--- a/src/dawn_native/Device.cpp
+++ b/src/dawn_native/Device.cpp
@@ -225,7 +225,7 @@
mDynamicUploader = std::make_unique<DynamicUploader>(this);
mCallbackTaskManager = std::make_unique<CallbackTaskManager>();
mDeprecationWarnings = std::make_unique<DeprecationWarnings>();
- mInternalPipelineStore = std::make_unique<InternalPipelineStore>();
+ mInternalPipelineStore = std::make_unique<InternalPipelineStore>(this);
mPersistentCache = std::make_unique<PersistentCache>(this);
ASSERT(GetPlatform() != nullptr);
diff --git a/src/dawn_native/EncodingContext.cpp b/src/dawn_native/EncodingContext.cpp
index 9e8812d..9e7b960 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 @@
}
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 @@
}
}
+ 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 @@
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 @@
// if Finish() has been called.
mCurrentEncoder = nullptr;
mTopLevelEncoder = nullptr;
+ CommitCommands(std::move(mPendingCommands));
if (mError != nullptr) {
return std::move(mError);
@@ -136,6 +170,12 @@
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 b97e317..522d29e 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 @@
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 @@
ComputePassUsages AcquireComputePassUsages();
private:
+ void CommitCommands(CommandAllocator allocator);
+
bool IsFinished() const;
void MoveToIterator();
@@ -104,7 +116,9 @@
ComputePassUsages mComputePassUsages;
bool mWereComputePassUsagesAcquired = false;
- CommandAllocator mAllocator;
+ CommandAllocator mPendingCommands;
+
+ std::vector<CommandAllocator> 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 0000000..235935f
--- /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 <algorithm>
+#include <utility>
+
+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::IndexedIndirectValidationBatch>&
+ 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 0000000..04c38e3
--- /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 <cstdint>
+#include <map>
+#include <set>
+#include <utility>
+#include <vector>
+
+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> bufferLocation;
+ };
+
+ struct IndexedIndirectValidationBatch {
+ uint64_t minOffset;
+ uint64_t maxOffset;
+ std::vector<IndexedIndirectDraw> 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<IndexedIndirectValidationBatch>& GetBatches() const;
+
+ private:
+ Ref<BufferBase> 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<IndexedIndirectValidationBatch> 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<BufferBase*, uint64_t>;
+ using IndexedIndirectBufferValidationInfoMap =
+ std::map<IndexedIndirectConfig, IndexedIndirectBufferValidationInfo>;
+
+ 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<RenderBundleBase*> 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 0000000..c3c2a04
--- /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 <cstdlib>
+#include <limits>
+
+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<u32>;
+ };
+
+ [[block]] struct IndirectParams {
+ data: array<u32>;
+ };
+
+ [[group(0), binding(0)]] var<storage, read> batch: BatchInfo;
+ [[group(0), binding(1)]] var<storage, read_write> clientParams: IndirectParams;
+ [[group(0), binding(2)]] var<storage, write> 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<u32>) {
+ 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<ComputePipelineBase*> 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<ChainedStruct*>(&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<BindGroupLayoutBase> bindGroupLayout;
+ DAWN_TRY_ASSIGN(bindGroupLayout,
+ device->CreateBindGroupLayout(&bindGroupLayoutDescriptor, true));
+
+ PipelineLayoutDescriptor pipelineDescriptor;
+ pipelineDescriptor.bindGroupLayoutCount = 1;
+ pipelineDescriptor.bindGroupLayouts = &bindGroupLayout.Get();
+ Ref<PipelineLayoutBase> 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<void, void (*)(void*)> batchData{nullptr, std::free};
+ std::vector<Batch> 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<Pass> 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<DeferredBufferLocationUpdate> 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<uint8_t*>(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<uint32_t>(batch.metadata->draws.size());
+
+ uint32_t* indirectOffsets = reinterpret_cast<uint32_t*>(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<uint32_t>(
+ (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<BindGroupLayoutBase> 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<const uint8_t*>(pass.batchData.get()),
+ pass.batchDataSize);
+
+ // TODO(dawn:723): change to not use AcquireRef for reentrant object creation.
+ ComputePassDescriptor descriptor = {};
+ Ref<ComputePassEncoder> 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<BindGroupBase> 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 0000000..bc62bf0
--- /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 0000000..edfd115
--- /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 <unordered_map>
+
+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 b3a7398..acf3b13 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 <unordered_map>
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<wgpu::TextureFormat, Ref<RenderPipelineBase>>
copyTextureForBrowserPipelines;
@@ -34,7 +42,18 @@
Ref<ShaderModuleBase> timestampCS;
Ref<ShaderModuleBase> 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<ComputePipelineBase> renderValidationPipeline;
+ Ref<ShaderModuleBase> renderValidationShader;
};
+
} // namespace dawn_native
#endif // DAWNNATIVE_INTERNALPIPELINESTORE_H_
diff --git a/src/dawn_native/RenderBundle.cpp b/src/dawn_native/RenderBundle.cpp
index 028dde7..b17031b 100644
--- a/src/dawn_native/RenderBundle.cpp
+++ b/src/dawn_native/RenderBundle.cpp
@@ -24,9 +24,11 @@
RenderBundleBase::RenderBundleBase(RenderBundleEncoder* encoder,
const RenderBundleDescriptor* descriptor,
Ref<AttachmentState> 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 @@
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 f971ed6..40db924 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 @@
RenderBundleBase(RenderBundleEncoder* encoder,
const RenderBundleDescriptor* descriptor,
Ref<AttachmentState> attachmentState,
- RenderPassResourceUsage resourceUsage);
+ RenderPassResourceUsage resourceUsage,
+ IndirectDrawMetadata indirectDrawMetadata);
static RenderBundleBase* MakeError(DeviceBase* device);
@@ -44,6 +46,7 @@
const AttachmentState* GetAttachmentState() const;
const RenderPassResourceUsage& GetResourceUsage() const;
+ const IndirectDrawMetadata& GetIndirectDrawMetadata();
protected:
~RenderBundleBase() override;
@@ -52,6 +55,7 @@
RenderBundleBase(DeviceBase* device, ErrorTag errorTag);
CommandIterator mCommands;
+ IndirectDrawMetadata mIndirectDrawMetadata;
Ref<AttachmentState> mAttachmentState;
RenderPassResourceUsage mResourceUsage;
};
diff --git a/src/dawn_native/RenderBundleEncoder.cpp b/src/dawn_native/RenderBundleEncoder.cpp
index daff3eb..7ddda31 100644
--- a/src/dawn_native/RenderBundleEncoder.cpp
+++ b/src/dawn_native/RenderBundleEncoder.cpp
@@ -130,7 +130,8 @@
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 06b0f9b..58849cf 100644
--- a/src/dawn_native/RenderEncoderBase.cpp
+++ b/src/dawn_native/RenderEncoderBase.cpp
@@ -157,16 +157,6 @@
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 @@
DrawIndexedIndirectCmd* cmd =
allocator->Allocate<DrawIndexedIndirectCmd>(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 4976ee2..30b7a3c 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 @@
CommandBufferStateTracker mCommandBufferState;
RenderPassResourceUsageTracker mUsageTracker;
+ IndirectDrawMetadata mIndirectDrawMetadata;
private:
Ref<AttachmentState> mAttachmentState;
diff --git a/src/dawn_native/RenderPassEncoder.cpp b/src/dawn_native/RenderPassEncoder.cpp
index b9cee4f..250d064 100644
--- a/src/dawn_native/RenderPassEncoder.cpp
+++ b/src/dawn_native/RenderPassEncoder.cpp
@@ -99,9 +99,11 @@
}
allocator->Allocate<EndRenderPassCmd>(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 @@
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 0000000..976214c
--- /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 0000000..7bb446d
--- /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 <cstdint>
+
+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<BufferBase> 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 175e3dd..4bdc3b0 100644
--- a/src/dawn_native/d3d12/CommandBufferD3D12.cpp
+++ b/src/dawn_native/d3d12/CommandBufferD3D12.cpp
@@ -981,6 +981,10 @@
break;
}
+ case Command::SetValidatedBufferLocationsInternal:
+ DoNextSetValidatedBufferLocationsInternal();
+ break;
+
case Command::WriteBuffer: {
WriteBufferCmd* write = mCommands.NextCommand<WriteBufferCmd>();
const uint64_t offset = write->offset;
diff --git a/src/dawn_native/metal/CommandBufferMTL.mm b/src/dawn_native/metal/CommandBufferMTL.mm
index 4dd47a5..eb9892b 100644
--- a/src/dawn_native/metal/CommandBufferMTL.mm
+++ b/src/dawn_native/metal/CommandBufferMTL.mm
@@ -987,6 +987,10 @@
break;
}
+ case Command::SetValidatedBufferLocationsInternal:
+ DoNextSetValidatedBufferLocationsInternal();
+ break;
+
case Command::WriteBuffer: {
WriteBufferCmd* write = mCommands.NextCommand<WriteBufferCmd>();
const uint64_t offset = write->offset;
diff --git a/src/dawn_native/opengl/CommandBufferGL.cpp b/src/dawn_native/opengl/CommandBufferGL.cpp
index 099d590..f4bb61f 100644
--- a/src/dawn_native/opengl/CommandBufferGL.cpp
+++ b/src/dawn_native/opengl/CommandBufferGL.cpp
@@ -843,6 +843,10 @@
break;
}
+ case Command::SetValidatedBufferLocationsInternal:
+ DoNextSetValidatedBufferLocationsInternal();
+ break;
+
case Command::WriteBuffer: {
WriteBufferCmd* write = mCommands.NextCommand<WriteBufferCmd>();
uint64_t offset = write->offset;
diff --git a/src/dawn_native/vulkan/CommandBufferVk.cpp b/src/dawn_native/vulkan/CommandBufferVk.cpp
index cfc2a71..36f6a39 100644
--- a/src/dawn_native/vulkan/CommandBufferVk.cpp
+++ b/src/dawn_native/vulkan/CommandBufferVk.cpp
@@ -824,6 +824,10 @@
break;
}
+ case Command::SetValidatedBufferLocationsInternal:
+ DoNextSetValidatedBufferLocationsInternal();
+ break;
+
case Command::WriteBuffer: {
WriteBufferCmd* write = mCommands.NextCommand<WriteBufferCmd>();
const uint64_t offset = write->offset;
diff --git a/src/tests/end2end/DrawIndexedIndirectTests.cpp b/src/tests/end2end/DrawIndexedIndirectTests.cpp
index 147d512..0a8a83f 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 @@
// 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<uint32_t>(
- 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<uint32_t> bufferList,
- uint64_t indexOffset,
- uint64_t indirectOffset,
- RGBA8 bottomLeftExpected,
- RGBA8 topRightExpected) {
- wgpu::Buffer indirectBuffer =
- utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Indirect, bufferList);
+ wgpu::Buffer CreateIndirectBuffer(std::initializer_list<uint32_t> indirectParamList) {
+ return utils::CreateBufferFromData<uint32_t>(
+ device, wgpu::BufferUsage::Indirect | wgpu::BufferUsage::Storage, indirectParamList);
+ }
+
+ wgpu::Buffer CreateIndexBuffer(std::initializer_list<uint32_t> indexList) {
+ return utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Index, indexList);
+ }
+
+ wgpu::CommandBuffer EncodeDrawCommands(std::initializer_list<uint32_t> bufferList,
+ wgpu::Buffer indexBuffer,
+ uint64_t indexOffset,
+ uint64_t indirectOffset) {
+ wgpu::Buffer indirectBuffer = CreateIndirectBuffer(bufferList);
wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
{
@@ -89,12 +91,28 @@
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<uint32_t> 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({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<wgpu::RenderBundle> 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<uniform> input: Input;
+ [[group(0), binding(1)]] var<storage, write> 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<uint32_t>(
+ 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 abac251..ffdb128 100644
--- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
+++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
@@ -28,87 +28,6 @@
}
};
-// 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<f32> {
- return vec4<f32>();
- })");
- 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.
