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>

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;

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

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,

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"

src/dawn_native/CommandBuffer.cpp

@@ -38,6 +38,14 @@ namespace dawn_native {
         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);

src/dawn_native/CommandBuffer.h

@@ -44,6 +44,8 @@ namespace dawn_native {
       protected:
         ~CommandBufferBase();
 
+        void DoNextSetValidatedBufferLocationsInternal();
+
         CommandIterator mCommands;
 
       private:

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

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

src/dawn_native/CommandEncoder.cpp

@@ -508,6 +508,7 @@ namespace dawn_native {
         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 @@ namespace dawn_native {
         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();

src/dawn_native/CommandEncoder.h

@@ -76,6 +76,9 @@ namespace dawn_native {
 
         CommandBufferBase* APIFinish(const CommandBufferDescriptor* descriptor = nullptr);
 
+        void EncodeSetValidatedBufferLocationsInternal(
+            std::vector<DeferredBufferLocationUpdate> updates);
+
       private:
         ResultOrError<Ref<CommandBufferBase>> FinishInternal(
             const CommandBufferDescriptor* descriptor);

src/dawn_native/Commands.cpp

@@ -158,6 +158,12 @@ namespace dawn_native {
                     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 @@ namespace dawn_native {
                 commands->NextCommand<SetStencilReferenceCmd>();
                 break;
 
+            case Command::SetValidatedBufferLocationsInternal:
+                commands->NextCommand<SetValidatedBufferLocationsInternalCmd>();
+                break;
+
             case Command::SetViewport:
                 commands->NextCommand<SetViewportCmd>();
                 break;

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<BufferLocation> location;
+        Ref<BufferBase> buffer;
+        uint64_t offset;
+    };
+
+    struct SetValidatedBufferLocationsInternalCmd {
+        std::vector<DeferredBufferLocationUpdate> updates;
+    };
+
     struct SetViewportCmd {
         float x, y, width, height, minDepth, maxDepth;
     };

src/dawn_native/ComputePassEncoder.cpp

@@ -67,7 +67,7 @@ namespace dawn_native {
 
                 return {};
             })) {
-            mEncodingContext->ExitPass(this, mUsageTracker.AcquireResourceUsage());
+            mEncodingContext->ExitComputePass(this, mUsageTracker.AcquireResourceUsage());
         }
     }
 

src/dawn_native/Device.cpp

@@ -225,7 +225,7 @@ namespace dawn_native {
         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);

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));
         }
 
-    void EncodingContext::ExitPass(const ObjectBase* passEncoder, ComputePassResourceUsage usages) {
+        mRenderPassUsages.push_back(usageTracker.AcquireResourceUsage());
+        return {};
+    }
+
+    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;
     }

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);
+        MaybeError ExitRenderPass(const ObjectBase* passEncoder,
-        void ExitPass(const ObjectBase* passEncoder, ComputePassResourceUsage usages);
+                                  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<CommandAllocator> mAllocators;
         CommandIterator mIterator;
         bool mWasMovedToIterator = false;
         bool mWereCommandsAcquired = false;

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

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_

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

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_

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

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 @@ namespace dawn_native {
         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_

src/dawn_native/RenderBundle.cpp

@@ -24,9 +24,11 @@ namespace dawn_native {
     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 @@ namespace dawn_native {
         return mResourceUsage;
     }
 
+    const IndirectDrawMetadata& RenderBundleBase::GetIndirectDrawMetadata() {
+        return mIndirectDrawMetadata;
+    }
+
 }  // namespace dawn_native

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> 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<AttachmentState> mAttachmentState;
         RenderPassResourceUsage mResourceUsage;
     };

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 {

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<DrawIndexedIndirectCmd>(Command::DrawIndexedIndirect);
+            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);
 

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<AttachmentState> mAttachmentState;

src/dawn_native/RenderPassEncoder.cpp

@@ -99,9 +99,11 @@ namespace dawn_native {
                 }
 
                 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 @@ namespace dawn_native {
                     mUsageTracker.AddRenderBundleTextureUsage(usages.textures[i],
                                                               usages.textureUsages[i]);
                 }
+
+                if (IsValidationEnabled()) {
+                    mIndirectDrawMetadata.AddBundle(renderBundles[i]);
+                }
             }
 
             return {};

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

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_

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<WriteBufferCmd>();
                     const uint64_t offset = write->offset;

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<WriteBufferCmd>();
                     const uint64_t offset = write->offset;

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<WriteBufferCmd>();
                     uint64_t offset = write->offset;

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<WriteBufferCmd>();
                     const uint64_t offset = write->offset;

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<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,
+    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,
+                                           uint64_t indirectOffset) {
-              RGBA8 bottomLeftExpected,
+        wgpu::Buffer indirectBuffer = CreateIndirectBuffer(bufferList);
-              RGBA8 topRightExpected) {
-        wgpu::Buffer indirectBuffer =
-            utils::CreateBufferFromData<uint32_t>(device, wgpu::BufferUsage::Indirect, 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<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_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<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(),

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<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.