D3D12: Duplicate first/baseVertex on Draw[Indexed]Indirect

Adds support for non-zero first/baseVertex on Draw[Indexed]Indirect by
duplicating the first/baseVertex indirect parameter onto a root
constant in the indirect buffer.

Change-Id: I280149065179806d3e57b07f1a396f9e2e4e8fcb
Bug: dawn:548
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/84240
Reviewed-by: Austin Eng <enga@chromium.org>
Commit-Queue: Enrico Galli <enrico.galli@intel.com>
Auto-Submit: Enrico Galli <enrico.galli@intel.com>

src/dawn/native/Commands.h

@@ -188,10 +188,7 @@ namespace dawn::native {
         uint64_t indirectOffset;
     };
 
-    struct DrawIndexedIndirectCmd {
+    struct DrawIndexedIndirectCmd : DrawIndirectCmd {};
-        Ref<BufferBase> indirectBuffer;
-        uint64_t indirectOffset;
-    };
 
     struct EndComputePassCmd {
         std::vector<TimestampWrite> timestampWrites;

src/dawn/native/Device.cpp

@@ -1830,4 +1830,13 @@ namespace dawn::native {
         return false;
     }
 
+    bool DeviceBase::MayRequireDuplicationOfIndirectParameters() const {
+        return false;
+    }
+
+    bool DeviceBase::ShouldDuplicateParametersForDrawIndirect(
+        const RenderPipelineBase* renderPipelineBase) const {
+        return false;
+    }
+
 }  // namespace dawn::native

src/dawn/native/Device.h

@@ -356,6 +356,11 @@ namespace dawn::native {
         virtual bool ShouldDuplicateNumWorkgroupsForDispatchIndirect(
             ComputePipelineBase* computePipeline) const;
 
+        virtual bool MayRequireDuplicationOfIndirectParameters() const;
+
+        virtual bool ShouldDuplicateParametersForDrawIndirect(
+            const RenderPipelineBase* renderPipelineBase) const;
+
         const CombinedLimits& GetLimits() const;
 
         AsyncTaskManager* GetAsyncTaskManager() const;

src/dawn/native/EncodingContext.cpp

@@ -91,12 +91,14 @@ namespace dawn::native {
 
     void EncodingContext::WillBeginRenderPass() {
         ASSERT(mCurrentEncoder == mTopLevelEncoder);
-        if (mDevice->IsValidationEnabled()) {
+        if (mDevice->IsValidationEnabled() ||
-            // When validation is enabled, we are going to want to capture all commands encoded
+            mDevice->MayRequireDuplicationOfIndirectParameters()) {
-            // between and including BeginRenderPassCmd and EndRenderPassCmd, and defer their
+            // When validation is enabled or indirect parameters require duplication, we are going
-            // sequencing util after we have a chance to insert any necessary validation
+            // to want to capture all commands encoded between and including BeginRenderPassCmd and
-            // commands. To support this we commit any current commands now, so that the
+            // EndRenderPassCmd, and defer their sequencing util after we have a chance to insert
-            // impending BeginRenderPassCmd starts in a fresh CommandAllocator.
+            // any necessary validation or duplication commands. To support this we commit any
+            // current commands now, so that the impending BeginRenderPassCmd starts in a fresh
+            // CommandAllocator.
             CommitCommands(std::move(mPendingCommands));
         }
     }
@@ -118,7 +120,8 @@ namespace dawn::native {
 
         mCurrentEncoder = mTopLevelEncoder;
 
-        if (mDevice->IsValidationEnabled()) {
+        if (mDevice->IsValidationEnabled() ||
+            mDevice->MayRequireDuplicationOfIndirectParameters()) {
             // 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

src/dawn/native/IndirectDrawMetadata.cpp

@@ -35,14 +35,14 @@ namespace dawn::native {
         : mIndirectBuffer(indirectBuffer) {
     }
 
-    void IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::AddIndexedIndirectDraw(
+    void IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::AddIndirectDraw(
         uint32_t maxDrawCallsPerIndirectValidationBatch,
         uint64_t maxBatchOffsetRange,
-        IndexedIndirectDraw draw) {
+        IndirectDraw draw) {
         const uint64_t newOffset = draw.clientBufferOffset;
         auto it = mBatches.begin();
         while (it != mBatches.end()) {
-            IndexedIndirectValidationBatch& batch = *it;
+            IndirectValidationBatch& batch = *it;
             if (batch.draws.size() >= maxDrawCallsPerIndirectValidationBatch) {
                 // This batch is full. If its minOffset is to the right of the new offset, we can
                 // just insert a new batch here.
@@ -82,7 +82,7 @@ namespace dawn::native {
             ++it;
         }
 
-        IndexedIndirectValidationBatch newBatch;
+        IndirectValidationBatch newBatch;
         newBatch.minOffset = newOffset;
         newBatch.maxOffset = newOffset;
         newBatch.draws.push_back(std::move(draw));
@@ -93,10 +93,10 @@ namespace dawn::native {
     void IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::AddBatch(
         uint32_t maxDrawCallsPerIndirectValidationBatch,
         uint64_t maxBatchOffsetRange,
-        const IndexedIndirectValidationBatch& newBatch) {
+        const IndirectValidationBatch& newBatch) {
         auto it = mBatches.begin();
         while (it != mBatches.end()) {
-            IndexedIndirectValidationBatch& batch = *it;
+            IndirectValidationBatch& batch = *it;
             uint64_t min = std::min(newBatch.minOffset, batch.minOffset);
             uint64_t max = std::max(newBatch.maxOffset, batch.maxOffset);
             if (max - min <= maxBatchOffsetRange && batch.draws.size() + newBatch.draws.size() <=
@@ -117,7 +117,7 @@ namespace dawn::native {
         mBatches.push_back(newBatch);
     }
 
-    const std::vector<IndirectDrawMetadata::IndexedIndirectValidationBatch>&
+    const std::vector<IndirectDrawMetadata::IndirectValidationBatch>&
     IndirectDrawMetadata::IndexedIndirectBufferValidationInfo::GetBatches() const {
         return mBatches;
     }
@@ -149,7 +149,7 @@ namespace dawn::native {
             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 : validationInfo.GetBatches()) {
+                for (const IndirectValidationBatch& batch : validationInfo.GetBatches()) {
                     it->second.AddBatch(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange, batch);
                 }
             } else {
@@ -162,6 +162,7 @@ namespace dawn::native {
                                                       uint64_t indexBufferSize,
                                                       BufferBase* indirectBuffer,
                                                       uint64_t indirectOffset,
+                                                      bool duplicateBaseVertexInstance,
                                                       DrawIndexedIndirectCmd* cmd) {
         uint64_t numIndexBufferElements;
         switch (indexFormat) {
@@ -175,7 +176,8 @@ namespace dawn::native {
                 UNREACHABLE();
         }
 
-        const IndexedIndirectConfig config(indirectBuffer, numIndexBufferElements);
+        const IndexedIndirectConfig config = {indirectBuffer, numIndexBufferElements,
+                                              duplicateBaseVertexInstance, DrawType::Indexed};
         auto it = mIndexedIndirectBufferValidationInfo.find(config);
         if (it == mIndexedIndirectBufferValidationInfo.end()) {
             auto result = mIndexedIndirectBufferValidationInfo.emplace(
@@ -183,11 +185,45 @@ namespace dawn::native {
             it = result.first;
         }
 
-        IndexedIndirectDraw draw;
+        IndirectDraw draw{};
         draw.clientBufferOffset = indirectOffset;
         draw.cmd = cmd;
-        it->second.AddIndexedIndirectDraw(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange,
+        it->second.AddIndirectDraw(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange, draw);
-                                          std::move(draw));
+    }
+
+    void IndirectDrawMetadata::AddIndirectDraw(BufferBase* indirectBuffer,
+                                               uint64_t indirectOffset,
+                                               bool duplicateBaseVertexInstance,
+                                               DrawIndirectCmd* cmd) {
+        const IndexedIndirectConfig config = {indirectBuffer, 0, duplicateBaseVertexInstance,
+                                              DrawType::NonIndexed};
+        auto it = mIndexedIndirectBufferValidationInfo.find(config);
+        if (it == mIndexedIndirectBufferValidationInfo.end()) {
+            auto result = mIndexedIndirectBufferValidationInfo.emplace(
+                config, IndexedIndirectBufferValidationInfo(indirectBuffer));
+            it = result.first;
+        }
+
+        IndirectDraw draw{};
+        draw.clientBufferOffset = indirectOffset;
+        draw.cmd = cmd;
+        it->second.AddIndirectDraw(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange, draw);
+    }
+
+    bool IndirectDrawMetadata::IndexedIndirectConfig::operator<(
+        const IndexedIndirectConfig& other) const {
+        return std::tie(clientIndirectBuffer, numIndexBufferElements, duplicateBaseVertexInstance,
+                        drawType) < std::tie(other.clientIndirectBuffer,
+                                             other.numIndexBufferElements,
+                                             other.duplicateBaseVertexInstance, other.drawType);
+    }
+
+    bool IndirectDrawMetadata::IndexedIndirectConfig::operator==(
+        const IndexedIndirectConfig& other) const {
+        return std::tie(clientIndirectBuffer, numIndexBufferElements, duplicateBaseVertexInstance,
+                        drawType) == std::tie(other.clientIndirectBuffer,
+                                              other.numIndexBufferElements,
+                                              other.duplicateBaseVertexInstance, other.drawType);
     }
 
 }  // namespace dawn::native

src/dawn/native/IndirectDrawMetadata.h

@@ -42,18 +42,18 @@ namespace dawn::native {
     // commands.
     class IndirectDrawMetadata : public NonCopyable {
       public:
-        struct IndexedIndirectDraw {
+        struct IndirectDraw {
             uint64_t clientBufferOffset;
             // This is a pointer to the command that should be populated with the validated
             // indirect scratch buffer. It is only valid up until the encoded command buffer
             // is submitted.
-            DrawIndexedIndirectCmd* cmd;
+            DrawIndirectCmd* cmd;
         };
 
-        struct IndexedIndirectValidationBatch {
+        struct IndirectValidationBatch {
             uint64_t minOffset;
             uint64_t maxOffset;
-            std::vector<IndexedIndirectDraw> draws;
+            std::vector<IndirectDraw> draws;
         };
 
         // Tracks information about every draw call in this render pass which uses the same indirect
@@ -65,18 +65,18 @@ namespace dawn::native {
 
             // 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(uint32_t maxDrawCallsPerIndirectValidationBatch,
+            void AddIndirectDraw(uint32_t maxDrawCallsPerIndirectValidationBatch,
                                  uint64_t maxBatchOffsetRange,
-                                        IndexedIndirectDraw draw);
+                                 IndirectDraw 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(uint32_t maxDrawCallsPerIndirectValidationBatch,
                           uint64_t maxBatchOffsetRange,
-                          const IndexedIndirectValidationBatch& batch);
+                          const IndirectValidationBatch& batch);
 
-            const std::vector<IndexedIndirectValidationBatch>& GetBatches() const;
+            const std::vector<IndirectValidationBatch>& GetBatches() const;
 
           private:
             Ref<BufferBase> mIndirectBuffer;
@@ -89,12 +89,23 @@ namespace dawn::native {
             // 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;
+            std::vector<IndirectValidationBatch> mBatches;
+        };
+
+        enum class DrawType {
+            NonIndexed,
+            Indexed,
+        };
+        struct IndexedIndirectConfig {
+            BufferBase* clientIndirectBuffer;
+            uint64_t numIndexBufferElements;
+            bool duplicateBaseVertexInstance;
+            DrawType drawType;
+
+            bool operator<(const IndexedIndirectConfig& other) const;
+            bool operator==(const IndexedIndirectConfig& other) const;
         };
 
-        // 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>;
 
@@ -111,8 +122,14 @@ namespace dawn::native {
                                     uint64_t indexBufferSize,
                                     BufferBase* indirectBuffer,
                                     uint64_t indirectOffset,
+                                    bool duplicateBaseVertexInstance,
                                     DrawIndexedIndirectCmd* cmd);
 
+        void AddIndirectDraw(BufferBase* indirectBuffer,
+                             uint64_t indirectOffset,
+                             bool duplicateBaseVertexInstance,
+                             DrawIndirectCmd* cmd);
+
       private:
         IndexedIndirectBufferValidationInfoMap mIndexedIndirectBufferValidationInfo;
         std::set<RenderBundleBase*> mAddedBundles;

src/dawn/native/IndirectDrawValidationEncoder.cpp

@@ -39,62 +39,93 @@ namespace dawn::native {
         // NOTE: This must match the workgroup_size attribute on the compute entry point below.
         constexpr uint64_t kWorkgroupSize = 64;
 
+        // Bitmasks for BatchInfo::flags
+        constexpr uint32_t kDuplicateBaseVertexInstance = 1;
+        constexpr uint32_t kIndexedDraw = 2;
+        constexpr uint32_t kValidationEnabled = 4;
+
         // Equivalent to the BatchInfo struct defined in the shader below.
         struct BatchInfo {
             uint64_t numIndexBufferElements;
             uint32_t numDraws;
-            uint32_t padding;
+            uint32_t flags;
         };
 
         // 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 kNumDrawIndirectParams = 4u;
 
             let kIndexCountEntry = 0u;
-            let kInstanceCountEntry = 1u;
             let kFirstIndexEntry = 2u;
-            let kBaseVertexEntry = 3u;
+
-            let kFirstInstanceEntry = 4u;
+            // Bitmasks for BatchInfo::flags
+            let kDuplicateBaseVertexInstance = 1u;
+            let kIndexedDraw = 2u;
+            let kValidationEnabled = 4u;
 
             struct BatchInfo {
-                numIndexBufferElementsLow: u32;
+                numIndexBufferElementsLow: u32,
-                numIndexBufferElementsHigh: u32;
+                numIndexBufferElementsHigh: u32,
-                numDraws: u32;
+                numDraws: u32,
-                padding: u32;
+                flags: u32,
-                indirectOffsets: array<u32>;
+                indirectOffsets: array<u32>,
-            };
+            }
 
             struct IndirectParams {
-                data: array<u32>;
+                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 numIndirectParamsPerDrawCallClient() -> u32 {
+                var numParams = kNumDrawIndirectParams;
+                // Indexed Draw has an extra parameter (firstIndex)
+                if (bool(batch.flags & kIndexedDraw)) {
+                    numParams = numParams + 1u;
+                }
+                return numParams;
+            }
+
+            fn numIndirectParamsPerDrawCallValidated() -> u32 {
+                var numParams = numIndirectParamsPerDrawCallClient();
+                // 2 extra parameter for duplicated first/baseVexter and firstInstance
+                if (bool(batch.flags & kDuplicateBaseVertexInstance)) {
+                    numParams = numParams + 2u;
+                }
+                return numParams;
+            }
+
             fn fail(drawIndex: u32) {
-                let index = drawIndex * kNumIndirectParamsPerDrawCall;
+                let numParams = numIndirectParamsPerDrawCallValidated();
-                validatedParams.data[index + kIndexCountEntry] = 0u;
+                let index = drawIndex * numParams;
-                validatedParams.data[index + kInstanceCountEntry] = 0u;
+                for(var i = 0u; i < numParams; i = i + 1u) {
-                validatedParams.data[index + kFirstIndexEntry] = 0u;
+                    validatedParams.data[index + i] = 0u;
-                validatedParams.data[index + kBaseVertexEntry] = 0u;
+                }
-                validatedParams.data[index + kFirstInstanceEntry] = 0u;
             }
 
             fn pass(drawIndex: u32) {
-                let vIndex = drawIndex * kNumIndirectParamsPerDrawCall;
+                let numClientParams = numIndirectParamsPerDrawCallClient();
+                var vIndex = drawIndex * numIndirectParamsPerDrawCallValidated();
                 let cIndex = batch.indirectOffsets[drawIndex];
-                validatedParams.data[vIndex + kIndexCountEntry] =
+
-                    clientParams.data[cIndex + kIndexCountEntry];
+                // The first 2 parameter is reserved for the duplicated first/baseVertex and firstInstance
-                validatedParams.data[vIndex + kInstanceCountEntry] =
+
-                    clientParams.data[cIndex + kInstanceCountEntry];
+                if (bool(batch.flags & kDuplicateBaseVertexInstance)) {
-                validatedParams.data[vIndex + kFirstIndexEntry] =
+                    // first/baseVertex and firstInstance are always last two parameters
-                    clientParams.data[cIndex + kFirstIndexEntry];
+                    let dupIndex = cIndex + numClientParams - 2u;
-                validatedParams.data[vIndex + kBaseVertexEntry] =
+                    validatedParams.data[vIndex] = clientParams.data[dupIndex];
-                    clientParams.data[cIndex + kBaseVertexEntry];
+                    validatedParams.data[vIndex + 1u] = clientParams.data[dupIndex + 1u];
-                validatedParams.data[vIndex + kFirstInstanceEntry] =
+
-                    clientParams.data[cIndex + kFirstInstanceEntry];
+                    vIndex = vIndex + 2u;
+                }
+
+                for(var i = 0u; i < numClientParams; i = i + 1u) {
+                    validatedParams.data[vIndex + i] = clientParams.data[cIndex + i];
+                }
             }
 
             @stage(compute) @workgroup_size(64, 1, 1)
@@ -103,13 +134,24 @@ namespace dawn::native {
                     return;
                 }
 
+                if(!bool(batch.flags & kValidationEnabled)) {
+                    pass(id.x);
+                    return;
+                }
+
                 let clientIndex = batch.indirectOffsets[id.x];
-                let firstInstance = clientParams.data[clientIndex + kFirstInstanceEntry];
+                // firstInstance is always the last parameter
+                let firstInstance = clientParams.data[clientIndex + numIndirectParamsPerDrawCallClient() - 1u];
                 if (firstInstance != 0u) {
                     fail(id.x);
                     return;
                 }
 
+                if (!bool(batch.flags & kIndexedDraw)) {
+                    pass(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.
@@ -198,7 +240,7 @@ namespace dawn::native {
                                                     RenderPassResourceUsageTracker* usageTracker,
                                                     IndirectDrawMetadata* indirectDrawMetadata) {
         struct Batch {
-            const IndirectDrawMetadata::IndexedIndirectValidationBatch* metadata;
+            const IndirectDrawMetadata::IndirectValidationBatch* metadata;
             uint64_t numIndexBufferElements;
             uint64_t dataBufferOffset;
             uint64_t dataSize;
@@ -210,6 +252,7 @@ namespace dawn::native {
         };
 
         struct Pass {
+            uint32_t flags;
             BufferBase* clientIndirectBuffer;
             uint64_t validatedParamsSize = 0;
             uint64_t batchDataSize = 0;
@@ -235,8 +278,17 @@ namespace dawn::native {
             device->GetLimits().v1.minStorageBufferOffsetAlignment;
 
         for (auto& [config, validationInfo] : bufferInfoMap) {
-            BufferBase* clientIndirectBuffer = config.first;
+            const uint64_t indirectDrawCommandSize =
-            for (const IndirectDrawMetadata::IndexedIndirectValidationBatch& batch :
+                config.drawType == IndirectDrawMetadata::DrawType::Indexed
+                    ? kDrawIndexedIndirectSize
+                    : kDrawIndirectSize;
+
+            uint64_t validatedIndirectSize = indirectDrawCommandSize;
+            if (config.duplicateBaseVertexInstance) {
+                validatedIndirectSize += 2 * sizeof(uint32_t);
+            }
+
+            for (const IndirectDrawMetadata::IndirectValidationBatch& batch :
                  validationInfo.GetBatches()) {
                 const uint64_t minOffsetFromAlignedBoundary =
                     batch.minOffset % minStorageBufferOffsetAlignment;
@@ -245,13 +297,13 @@ namespace dawn::native {
 
                 Batch newBatch;
                 newBatch.metadata = &batch;
-                newBatch.numIndexBufferElements = config.second;
+                newBatch.numIndexBufferElements = config.numIndexBufferElements;
                 newBatch.dataSize = GetBatchDataSize(batch.draws.size());
                 newBatch.clientIndirectOffset = minOffsetAlignedDown;
                 newBatch.clientIndirectSize =
-                    batch.maxOffset + kDrawIndexedIndirectSize - minOffsetAlignedDown;
+                    batch.maxOffset + indirectDrawCommandSize - minOffsetAlignedDown;
 
-                newBatch.validatedParamsSize = batch.draws.size() * kDrawIndexedIndirectSize;
+                newBatch.validatedParamsSize = batch.draws.size() * validatedIndirectSize;
                 newBatch.validatedParamsOffset =
                     Align(validatedParamsSize, minStorageBufferOffsetAlignment);
                 validatedParamsSize = newBatch.validatedParamsOffset + newBatch.validatedParamsSize;
@@ -260,7 +312,8 @@ namespace dawn::native {
                 }
 
                 Pass* currentPass = passes.empty() ? nullptr : &passes.back();
-                if (currentPass && currentPass->clientIndirectBuffer == clientIndirectBuffer) {
+                if (currentPass &&
+                    currentPass->clientIndirectBuffer == config.clientIndirectBuffer) {
                     uint64_t nextBatchDataOffset =
                         Align(currentPass->batchDataSize, minStorageBufferOffsetAlignment);
                     uint64_t newPassBatchDataSize = nextBatchDataOffset + newBatch.dataSize;
@@ -276,10 +329,20 @@ namespace dawn::native {
                 // We need to start a new pass for this batch.
                 newBatch.dataBufferOffset = 0;
 
-                Pass newPass;
+                Pass newPass{};
-                newPass.clientIndirectBuffer = clientIndirectBuffer;
+                newPass.clientIndirectBuffer = config.clientIndirectBuffer;
                 newPass.batchDataSize = newBatch.dataSize;
                 newPass.batches.push_back(newBatch);
+                newPass.flags = 0;
+                if (config.duplicateBaseVertexInstance) {
+                    newPass.flags |= kDuplicateBaseVertexInstance;
+                }
+                if (config.drawType == IndirectDrawMetadata::DrawType::Indexed) {
+                    newPass.flags |= kIndexedDraw;
+                }
+                if (device->IsValidationEnabled()) {
+                    newPass.flags |= kValidationEnabled;
+                }
                 passes.push_back(std::move(newPass));
             }
         }
@@ -308,6 +371,7 @@ namespace dawn::native {
                 batch.batchInfo = new (&batchData[batch.dataBufferOffset]) BatchInfo();
                 batch.batchInfo->numIndexBufferElements = batch.numIndexBufferElements;
                 batch.batchInfo->numDraws = static_cast<uint32_t>(batch.metadata->draws.size());
+                batch.batchInfo->flags = pass.flags;
 
                 uint32_t* indirectOffsets = reinterpret_cast<uint32_t*>(batch.batchInfo + 1);
                 uint64_t validatedParamsOffset = batch.validatedParamsOffset;
@@ -318,8 +382,11 @@ namespace dawn::native {
 
                     draw.cmd->indirectBuffer = validatedParamsBuffer.GetBuffer();
                     draw.cmd->indirectOffset = validatedParamsOffset;
-
+                    if (pass.flags & kIndexedDraw) {
                         validatedParamsOffset += kDrawIndexedIndirectSize;
+                    } else {
+                        validatedParamsOffset += kDrawIndirectSize;
+                    }
                 }
             }
         }
@@ -347,9 +414,10 @@ namespace dawn::native {
         bindGroupDescriptor.entryCount = 3;
         bindGroupDescriptor.entries = bindings;
 
-        // Finally, we can now encode our validation passes. Each pass first does a single
+        // Finally, we can now encode our validation and duplication passes. Each pass first does a
-        // WriteBuffer to get batch data over to the GPU, followed by a single compute pass. The
+        // two WriteBuffer to get batch and pass data over to the GPU, followed by a single compute
-        // compute pass encodes a separate SetBindGroup and Dispatch command for each batch.
+        // pass. The compute pass encodes a separate SetBindGroup and Dispatch command for each
+        // batch.
         for (const Pass& pass : passes) {
             commandEncoder->APIWriteBuffer(batchDataBuffer.GetBuffer(), 0,
                                            static_cast<const uint8_t*>(pass.batchData.get()),

src/dawn/native/RenderEncoderBase.cpp

@@ -172,9 +172,28 @@ namespace dawn::native {
                 }
 
                 DrawIndirectCmd* cmd = allocator->Allocate<DrawIndirectCmd>(Command::DrawIndirect);
+
+                bool duplicateBaseVertexInstance =
+                    GetDevice()->ShouldDuplicateParametersForDrawIndirect(
+                        mCommandBufferState.GetRenderPipeline());
+                if (IsValidationEnabled() || duplicateBaseVertexInstance) {
+                    // Later, EncodeIndirectDrawValidationCommands will allocate a scratch storage
+                    // buffer which will store the validated or duplicated indirect data. The buffer
+                    // and offset will be updated to point to it.
+                    // |EncodeIndirectDrawValidationCommands| is called at the end of encoding the
+                    // render pass, while the |cmd| pointer is still valid.
+                    cmd->indirectBuffer = nullptr;
+
+                    mIndirectDrawMetadata.AddIndirectDraw(indirectBuffer, indirectOffset,
+                                                          duplicateBaseVertexInstance, cmd);
+                } else {
                     cmd->indirectBuffer = indirectBuffer;
                     cmd->indirectOffset = indirectOffset;
+                }
 
+                // TODO(crbug.com/dawn/1166): Adding the indirectBuffer is needed for correct usage
+                // validation, but it will unnecessarily transition to indirectBuffer usage in the
+                // backend.
                 mUsageTracker.BufferUsedAs(indirectBuffer, wgpu::BufferUsage::Indirect);
 
                 return {};
@@ -204,10 +223,14 @@ namespace dawn::native {
 
                 DrawIndexedIndirectCmd* cmd =
                     allocator->Allocate<DrawIndexedIndirectCmd>(Command::DrawIndexedIndirect);
-                if (IsValidationEnabled()) {
+
+                bool duplicateBaseVertexInstance =
+                    GetDevice()->ShouldDuplicateParametersForDrawIndirect(
+                        mCommandBufferState.GetRenderPipeline());
+                if (IsValidationEnabled() || duplicateBaseVertexInstance) {
                     // Later, EncodeIndirectDrawValidationCommands will allocate a scratch storage
-                    // buffer which will store the validated indirect data. The buffer and offset
+                    // buffer which will store the validated or duplicated indirect data. The buffer
-                    // will be updated to point to it.
+                    // and offset will be updated to point to it.
                     // |EncodeIndirectDrawValidationCommands| is called at the end of encoding the
                     // render pass, while the |cmd| pointer is still valid.
                     cmd->indirectBuffer = nullptr;
@@ -215,7 +238,7 @@ namespace dawn::native {
                     mIndirectDrawMetadata.AddIndexedIndirectDraw(
                         mCommandBufferState.GetIndexFormat(),
                         mCommandBufferState.GetIndexBufferSize(), indirectBuffer, indirectOffset,
-                        cmd);
+                        duplicateBaseVertexInstance, cmd);
                 } else {
                     cmd->indirectBuffer = indirectBuffer;
                     cmd->indirectOffset = indirectOffset;

src/dawn/native/d3d12/CommandBufferD3D12.cpp

@@ -149,19 +149,10 @@ namespace dawn::native::d3d12 {
             if (!firstOffsetInfo.usesVertexIndex && !firstOffsetInfo.usesInstanceIndex) {
                 return;
             }
-            std::array<uint32_t, 2> offsets{};
+            std::array<uint32_t, 2> offsets{firstVertex, firstInstance};
-            uint32_t count = 0;
-            if (firstOffsetInfo.usesVertexIndex) {
-                offsets[firstOffsetInfo.vertexIndexOffset / sizeof(uint32_t)] = firstVertex;
-                ++count;
-            }
-            if (firstOffsetInfo.usesInstanceIndex) {
-                offsets[firstOffsetInfo.instanceIndexOffset / sizeof(uint32_t)] = firstInstance;
-                ++count;
-            }
             PipelineLayout* layout = ToBackend(pipeline->GetLayout());
             commandList->SetGraphicsRoot32BitConstants(layout->GetFirstIndexOffsetParameterIndex(),
-                                                       count, offsets.data(), 0);
+                                                       offsets.size(), offsets.data(), 0);
         }
 
         bool ShouldCopyUsingTemporaryBuffer(DeviceBase* device,
@@ -1451,13 +1442,9 @@ namespace dawn::native::d3d12 {
                     DAWN_TRY(bindingTracker->Apply(commandContext));
                     vertexBufferTracker.Apply(commandList, lastPipeline);
 
-                    // TODO(dawn:548): remove this once builtins are emulated for indirect draws.
-                    // Zero the index offset values to avoid reusing values from the previous draw
-                    RecordFirstIndexOffset(commandList, lastPipeline, 0, 0);
-
                     Buffer* buffer = ToBackend(draw->indirectBuffer.Get());
                     ComPtr<ID3D12CommandSignature> signature =
-                        ToBackend(GetDevice())->GetDrawIndirectSignature();
+                        lastPipeline->GetDrawIndirectCommandSignature();
                     commandList->ExecuteIndirect(signature.Get(), 1, buffer->GetD3D12Resource(),
                                                  draw->indirectOffset, nullptr, 0);
                     break;
@@ -1469,15 +1456,11 @@ namespace dawn::native::d3d12 {
                     DAWN_TRY(bindingTracker->Apply(commandContext));
                     vertexBufferTracker.Apply(commandList, lastPipeline);
 
-                    // TODO(dawn:548): remove this once builtins are emulated for indirect draws.
-                    // Zero the index offset values to avoid reusing values from the previous draw
-                    RecordFirstIndexOffset(commandList, lastPipeline, 0, 0);
-
                     Buffer* buffer = ToBackend(draw->indirectBuffer.Get());
                     ASSERT(buffer != nullptr);
 
                     ComPtr<ID3D12CommandSignature> signature =
-                        ToBackend(GetDevice())->GetDrawIndexedIndirectSignature();
+                        lastPipeline->GetDrawIndexedIndirectCommandSignature();
                     commandList->ExecuteIndirect(signature.Get(), 1, buffer->GetD3D12Resource(),
                                                  draw->indirectOffset, nullptr, 0);
                     break;

src/dawn/native/d3d12/DeviceD3D12.cpp

@@ -760,4 +760,14 @@ namespace dawn::native::d3d12 {
         SetDebugName(this, mD3d12Device.Get(), "Dawn_Device", GetLabel());
     }
 
+    bool Device::MayRequireDuplicationOfIndirectParameters() const {
+        return true;
+    }
+
+    bool Device::ShouldDuplicateParametersForDrawIndirect(
+        const RenderPipelineBase* renderPipelineBase) const {
+        return ToBackend(renderPipelineBase)->GetFirstOffsetInfo().usesVertexIndex ||
+               ToBackend(renderPipelineBase)->GetFirstOffsetInfo().usesInstanceIndex;
+    }
+
 }  // namespace dawn::native::d3d12

src/dawn/native/d3d12/DeviceD3D12.h

@@ -149,7 +149,12 @@ namespace dawn::native::d3d12 {
         bool ShouldDuplicateNumWorkgroupsForDispatchIndirect(
             ComputePipelineBase* computePipeline) const override;
 
-        // Dawn API
+        bool MayRequireDuplicationOfIndirectParameters() const override;
+
+        bool ShouldDuplicateParametersForDrawIndirect(
+            const RenderPipelineBase* renderPipelineBase) const override;
+
+        // Dawn APIs
         void SetLabelImpl() override;
 
       private:

src/dawn/native/d3d12/PipelineLayoutD3D12.cpp

@@ -376,4 +376,70 @@ namespace dawn::native::d3d12 {
         return mDispatchIndirectCommandSignatureWithNumWorkgroups.Get();
     }
 
+    ID3D12CommandSignature*
+    PipelineLayout::GetDrawIndirectCommandSignatureWithInstanceVertexOffsets() {
+        // mDrawIndirectCommandSignatureWithInstanceVertexOffsets won't be created until it is
+        // needed.
+        if (mDrawIndirectCommandSignatureWithInstanceVertexOffsets.Get() != nullptr) {
+            return mDrawIndirectCommandSignatureWithInstanceVertexOffsets.Get();
+        }
+
+        D3D12_INDIRECT_ARGUMENT_DESC argumentDescs[2] = {};
+        argumentDescs[0].Type = D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT;
+        argumentDescs[0].Constant.RootParameterIndex = GetFirstIndexOffsetParameterIndex();
+        argumentDescs[0].Constant.Num32BitValuesToSet = 2;
+        argumentDescs[0].Constant.DestOffsetIn32BitValues = 0;
+
+        // A command signature must contain exactly 1 Draw / Dispatch / DispatchMesh / DispatchRays
+        // command. That command must come last.
+        argumentDescs[1].Type = D3D12_INDIRECT_ARGUMENT_TYPE_DRAW;
+
+        D3D12_COMMAND_SIGNATURE_DESC programDesc = {};
+        programDesc.ByteStride = 6 * sizeof(uint32_t);
+        programDesc.NumArgumentDescs = 2;
+        programDesc.pArgumentDescs = argumentDescs;
+
+        // The root signature must be specified if and only if the command signature changes one of
+        // the root arguments.
+        ToBackend(GetDevice())
+            ->GetD3D12Device()
+            ->CreateCommandSignature(
+                &programDesc, GetRootSignature(),
+                IID_PPV_ARGS(&mDrawIndirectCommandSignatureWithInstanceVertexOffsets));
+        return mDrawIndirectCommandSignatureWithInstanceVertexOffsets.Get();
+    }
+
+    ID3D12CommandSignature*
+    PipelineLayout::GetDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets() {
+        // mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets won't be created until it
+        // is needed.
+        if (mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets.Get() != nullptr) {
+            return mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets.Get();
+        }
+
+        D3D12_INDIRECT_ARGUMENT_DESC argumentDescs[2] = {};
+        argumentDescs[0].Type = D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT;
+        argumentDescs[0].Constant.RootParameterIndex = GetFirstIndexOffsetParameterIndex();
+        argumentDescs[0].Constant.Num32BitValuesToSet = 2;
+        argumentDescs[0].Constant.DestOffsetIn32BitValues = 0;
+
+        // A command signature must contain exactly 1 Draw / Dispatch / DispatchMesh / DispatchRays
+        // command. That command must come last.
+        argumentDescs[1].Type = D3D12_INDIRECT_ARGUMENT_TYPE_DRAW_INDEXED;
+
+        D3D12_COMMAND_SIGNATURE_DESC programDesc = {};
+        programDesc.ByteStride = 7 * sizeof(uint32_t);
+        programDesc.NumArgumentDescs = 2;
+        programDesc.pArgumentDescs = argumentDescs;
+
+        // The root signature must be specified if and only if the command signature changes one of
+        // the root arguments.
+        ToBackend(GetDevice())
+            ->GetD3D12Device()
+            ->CreateCommandSignature(
+                &programDesc, GetRootSignature(),
+                IID_PPV_ARGS(&mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets));
+        return mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets.Get();
+    }
+
 }  // namespace dawn::native::d3d12

src/dawn/native/d3d12/PipelineLayoutD3D12.h

@@ -56,6 +56,10 @@ namespace dawn::native::d3d12 {
 
         ID3D12CommandSignature* GetDispatchIndirectCommandSignatureWithNumWorkgroups();
 
+        ID3D12CommandSignature* GetDrawIndirectCommandSignatureWithInstanceVertexOffsets();
+
+        ID3D12CommandSignature* GetDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets();
+
         struct PerBindGroupDynamicStorageBufferLengthInfo {
             // First register offset for a bind group's dynamic storage buffer lengths.
             // This is the index into the array of root constants where this bind group's
@@ -95,6 +99,9 @@ namespace dawn::native::d3d12 {
         uint32_t mDynamicStorageBufferLengthsParameterIndex;
         ComPtr<ID3D12RootSignature> mRootSignature;
         ComPtr<ID3D12CommandSignature> mDispatchIndirectCommandSignatureWithNumWorkgroups;
+        ComPtr<ID3D12CommandSignature> mDrawIndirectCommandSignatureWithInstanceVertexOffsets;
+        ComPtr<ID3D12CommandSignature>
+            mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets;
     };
 
 }  // namespace dawn::native::d3d12

src/dawn/native/d3d12/RenderPipelineD3D12.cpp

@@ -463,6 +463,24 @@ namespace dawn::native::d3d12 {
         SetDebugName(ToBackend(GetDevice()), GetPipelineState(), "Dawn_RenderPipeline", GetLabel());
     }
 
+    ComPtr<ID3D12CommandSignature> RenderPipeline::GetDrawIndirectCommandSignature() {
+        if (mFirstOffsetInfo.usesVertexIndex || mFirstOffsetInfo.usesInstanceIndex) {
+            return ToBackend(GetLayout())
+                ->GetDrawIndirectCommandSignatureWithInstanceVertexOffsets();
+        }
+
+        return ToBackend(GetDevice())->GetDrawIndirectSignature();
+    }
+
+    ComPtr<ID3D12CommandSignature> RenderPipeline::GetDrawIndexedIndirectCommandSignature() {
+        if (mFirstOffsetInfo.usesVertexIndex || mFirstOffsetInfo.usesInstanceIndex) {
+            return ToBackend(GetLayout())
+                ->GetDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets();
+        }
+
+        return ToBackend(GetDevice())->GetDrawIndexedIndirectSignature();
+    }
+
     D3D12_INPUT_LAYOUT_DESC RenderPipeline::ComputeInputLayout(
         std::array<D3D12_INPUT_ELEMENT_DESC, kMaxVertexAttributes>* inputElementDescriptors) {
         unsigned int count = 0;

src/dawn/native/d3d12/RenderPipelineD3D12.h

@@ -43,6 +43,10 @@ namespace dawn::native::d3d12 {
         // Dawn API
         void SetLabelImpl() override;
 
+        ComPtr<ID3D12CommandSignature> GetDrawIndirectCommandSignature();
+
+        ComPtr<ID3D12CommandSignature> GetDrawIndexedIndirectCommandSignature();
+
       private:
         ~RenderPipeline() override;
 

src/dawn/tests/end2end/BufferZeroInitTests.cpp

@@ -1256,6 +1256,10 @@ TEST_P(BufferZeroInitTest, SetIndexBuffer) {
 // Test the buffer will be lazily initialized correctly when its first use is an indirect buffer for
 // DrawIndirect.
 TEST_P(BufferZeroInitTest, IndirectBufferForDrawIndirect) {
+    // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
+    // the offset= that Tint/GLSL produces.
+    DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
+
     // Bind the whole buffer as an indirect buffer.
     {
         constexpr uint64_t kOffset = 0u;
@@ -1274,7 +1278,7 @@ TEST_P(BufferZeroInitTest, IndirectBufferForDrawIndirect) {
 TEST_P(BufferZeroInitTest, IndirectBufferForDrawIndexedIndirect) {
     // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
     // the offset= that Tint/GLSL produces.
-    DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL());
+    DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
 
     // Bind the whole buffer as an indirect buffer.
     {

src/dawn/tests/end2end/DrawIndirectTests.cpp

@@ -89,6 +89,10 @@ class DrawIndirectTest : public DawnTest {
 
 // The basic triangle draw.
 TEST_P(DrawIndirectTest, Uint32) {
+    // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
+    // the offsets that Tint/GLSL produces.
+    DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
+
     RGBA8 filled(0, 255, 0, 255);
     RGBA8 notFilled(0, 0, 0, 0);
 
@@ -106,6 +110,10 @@ TEST_P(DrawIndirectTest, Uint32) {
 }
 
 TEST_P(DrawIndirectTest, IndirectOffset) {
+    // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
+    // the offsets that Tint/GLSL produces.
+    DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
+
     RGBA8 filled(0, 255, 0, 255);
     RGBA8 notFilled(0, 0, 0, 0);
 

src/dawn/tests/end2end/FirstIndexOffsetTests.cpp

@@ -35,6 +35,10 @@ enum class CheckIndex : uint32_t {
     Instance = 0x0000002,
 };
 
+bool IsIndirectDraw(DrawMode mode) {
+    return mode == DrawMode::NonIndexedIndirect || mode == DrawMode::IndexedIndirect;
+}
+
 namespace dawn {
     template <>
     struct IsDawnBitmask<CheckIndex> {
@@ -51,6 +55,10 @@ class FirstIndexOffsetTests : public DawnTest {
   protected:
     void SetUp() override {
         DawnTest::SetUp();
+        // TODO(crbug.com/dawn/1292): Some Intel OpenGL drivers don't seem to like
+        // the offsets that Tint/GLSL produces.
+        DAWN_SUPPRESS_TEST_IF(IsIntel() && IsOpenGL() && IsLinux());
+
         // TODO(tint:451): Remove once "flat" is supported under OpenGL(ES).
         DAWN_SUPPRESS_TEST_IF(IsOpenGL() || IsOpenGLES());
     }
@@ -180,14 +188,10 @@ struct FragInputs {
         case DrawMode::Indexed:
             break;
         case DrawMode::NonIndexedIndirect:
-            // With DrawIndirect firstInstance is reserved and must be 0 according to spec.
-            ASSERT_EQ(firstInstance, 0u);
             indirectBuffer = utils::CreateBufferFromData<uint32_t>(
                 device, wgpu::BufferUsage::Indirect, {1, 1, firstVertex, firstInstance});
             break;
         case DrawMode::IndexedIndirect:
-            // With DrawIndexedIndirect firstInstance is reserved and must be 0 according to spec.
-            ASSERT_EQ(firstInstance, 0u);
             indirectBuffer = utils::CreateBufferFromData<uint32_t>(
                 device, wgpu::BufferUsage::Indirect, {1, 1, 0, firstVertex, firstInstance});
             break;
@@ -205,7 +209,8 @@ struct FragInputs {
     pass.SetBindGroup(0, bindGroup);
     // Do a first draw to make sure the offset values are correctly updated on the next draw.
     // We should only see the values from the second draw.
-    pass.Draw(1, 1, firstVertex + 1, firstInstance + 1);
+    std::array<uint32_t, 2> firstDrawValues = {firstVertex + 1, firstInstance + 1};
+    pass.Draw(1, 1, firstDrawValues[0], firstDrawValues[1]);
     switch (mode) {
         case DrawMode::NonIndexed:
             pass.Draw(1, 1, firstVertex, firstInstance);
@@ -229,11 +234,16 @@ struct FragInputs {
     queue.Submit(1, &commands);
 
     std::array<uint32_t, 2> expected = {firstVertex, firstInstance};
-    // TODO(dawn:548): remove this once builtins are emulated for indirect draws.
+
-    // Until then the expected values should always be {0, 0}.
+    // Per the specification, if validation is enabled and indirect-first-instance is not enabled,
-    if (IsD3D12() && (mode == DrawMode::NonIndexedIndirect || mode == DrawMode::IndexedIndirect)) {
+    // Draw[Indexed]Indirect with firstInstance > 0 will be a no-op. The buffer should still have
-        expected = {0, 0};
+    // the values from the first draw.
+    if (firstInstance > 0 && IsIndirectDraw(mode) &&
+        !device.HasFeature(wgpu::FeatureName::IndirectFirstInstance) &&
+        !HasToggleEnabled("skip_validation")) {
+        expected = {checkIndex & CheckIndex::Vertex ? firstDrawValues[0] : 0, firstDrawValues[1]};
     }
+
     EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, expected.size());
 }
 
@@ -268,18 +278,38 @@ TEST_P(FirstIndexOffsetTests, IndexedBothOffset) {
     TestBothIndices(DrawMode::Indexed, 7, 11);
 }
 
-// There are no instance_index tests because the spec forces it to be 0.
-
 // Test that vertex_index starts at 7 when drawn using DrawIndirect()
 TEST_P(FirstIndexOffsetTests, NonIndexedIndirectVertexOffset) {
     TestVertexIndex(DrawMode::NonIndexedIndirect, 7);
 }
 
+// Test that instance_index starts at 11 when drawn using DrawIndirect()
+TEST_P(FirstIndexOffsetTests, NonIndexedIndirectInstanceOffset) {
+    TestInstanceIndex(DrawMode::NonIndexedIndirect, 11);
+}
+
+// Test that vertex_index and instance_index start at 7 and 11 respectively when drawn using
+// DrawIndirect()
+TEST_P(FirstIndexOffsetTests, NonIndexedIndirectBothOffset) {
+    TestBothIndices(DrawMode::NonIndexedIndirect, 7, 11);
+}
+
 // Test that vertex_index starts at 7 when drawn using DrawIndexedIndirect()
 TEST_P(FirstIndexOffsetTests, IndexedIndirectVertex) {
     TestVertexIndex(DrawMode::IndexedIndirect, 7);
 }
 
+// Test that instance_index starts at 11 when drawn using DrawIndexed()
+TEST_P(FirstIndexOffsetTests, IndexedIndirectInstance) {
+    TestInstanceIndex(DrawMode::IndexedIndirect, 11);
+}
+
+// Test that vertex_index and instance_index start at 7 and 11 respectively when drawn using
+// DrawIndexed()
+TEST_P(FirstIndexOffsetTests, IndexedIndirectBothOffset) {
+    TestBothIndices(DrawMode::IndexedIndirect, 7, 11);
+}
+
 DAWN_INSTANTIATE_TEST(FirstIndexOffsetTests,
                       D3D12Backend(),
                       MetalBackend(),

src/tint/transform/first_index_offset.cc

@@ -128,18 +128,11 @@ void FirstIndexOffset::Run(CloneContext& ctx,
 
   if (has_vertex_index || has_instance_index) {
     // Add uniform buffer members and calculate byte offsets
-    uint32_t offset = 0;
     ast::StructMemberList members;
-    if (has_vertex_index) {
     members.push_back(ctx.dst->Member(kFirstVertexName, ctx.dst->ty.u32()));
-      vertex_index_offset = offset;
+    vertex_index_offset = 0;
-      offset += 4;
-    }
-    if (has_instance_index) {
     members.push_back(ctx.dst->Member(kFirstInstanceName, ctx.dst->ty.u32()));
-      instance_index_offset = offset;
+    instance_index_offset = 4;
-      offset += 4;
-    }
     auto* struct_ = ctx.dst->Structure(ctx.dst->Sym(), std::move(members));
 
     // Create a global to hold the uniform buffer

webgpu-cts/expectations.txt

@@ -304,51 +304,6 @@ crbug.com/dawn/1345 webgpu:api,validation,createComputePipeline:entry_point_name
 crbug.com/dawn/1345 webgpu:api,validation,createComputePipeline:entry_point_name_must_match:isAsync=false;shaderModuleEntryPoint="s%C3%A9quen%C3%A7age";* [ Failure ]
 crbug.com/dawn/1345 webgpu:api,validation,createComputePipeline:entry_point_name_must_match:isAsync=true;shaderModuleEntryPoint="s%C3%A9quen%C3%A7age";* [ Failure ]
 
-################################################################################
-# windows draw failures
-# KEEP
-################################################################################
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=0;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=false;indirect=true;vertex_buffer_offset=0;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=false;indirect=true;vertex_buffer_offset=32;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=false;indirect=true;vertex_buffer_offset=0;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=false;indirect=true;vertex_buffer_offset=32;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=3;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=false;indirect=true;vertex_buffer_offset=0;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=false;indirect=true;vertex_buffer_offset=32;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=1;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=false;indirect=true;vertex_buffer_offset=0;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=false;indirect=true;vertex_buffer_offset=32;index_buffer_offset="_undef_";base_vertex="_undef_" [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=0;index_buffer_offset=16;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=0;base_vertex=9 [ Failure ]
-crbug.com/dawn/548 [ win ] webgpu:api,operation,rendering,draw:arguments:first=3;count=6;first_instance=0;instance_count=4;indexed=true;indirect=true;vertex_buffer_offset=32;index_buffer_offset=16;base_vertex=9 [ Failure ]
-
 ################################################################################
 # external_texture failures
 # KEEP