From 45ec7c352815154142de230588a769537da09389 Mon Sep 17 00:00:00 2001 From: Enrico Galli Date: Thu, 21 Apr 2022 02:25:35 +0000 Subject: [PATCH] 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 Commit-Queue: Enrico Galli Auto-Submit: Enrico Galli --- src/dawn/native/Commands.h | 5 +- src/dawn/native/Device.cpp | 9 + src/dawn/native/Device.h | 5 + src/dawn/native/EncodingContext.cpp | 17 +- src/dawn/native/IndirectDrawMetadata.cpp | 60 +++++-- src/dawn/native/IndirectDrawMetadata.h | 43 +++-- .../native/IndirectDrawValidationEncoder.cpp | 158 +++++++++++++----- src/dawn/native/RenderEncoderBase.cpp | 35 +++- src/dawn/native/d3d12/CommandBufferD3D12.cpp | 25 +-- src/dawn/native/d3d12/DeviceD3D12.cpp | 10 ++ src/dawn/native/d3d12/DeviceD3D12.h | 7 +- src/dawn/native/d3d12/PipelineLayoutD3D12.cpp | 66 ++++++++ src/dawn/native/d3d12/PipelineLayoutD3D12.h | 7 + src/dawn/native/d3d12/RenderPipelineD3D12.cpp | 18 ++ src/dawn/native/d3d12/RenderPipelineD3D12.h | 4 + .../tests/end2end/BufferZeroInitTests.cpp | 6 +- src/dawn/tests/end2end/DrawIndirectTests.cpp | 8 + .../tests/end2end/FirstIndexOffsetTests.cpp | 52 ++++-- src/tint/transform/first_index_offset.cc | 15 +- webgpu-cts/expectations.txt | 45 ----- 20 files changed, 418 insertions(+), 177 deletions(-) diff --git a/src/dawn/native/Commands.h b/src/dawn/native/Commands.h index a75a407d5b..7be232d492 100644 --- a/src/dawn/native/Commands.h +++ b/src/dawn/native/Commands.h @@ -188,10 +188,7 @@ namespace dawn::native { uint64_t indirectOffset; }; - struct DrawIndexedIndirectCmd { - Ref indirectBuffer; - uint64_t indirectOffset; - }; + struct DrawIndexedIndirectCmd : DrawIndirectCmd {}; struct EndComputePassCmd { std::vector timestampWrites; diff --git a/src/dawn/native/Device.cpp b/src/dawn/native/Device.cpp index f48efc6a05..4f5d84e914 100644 --- a/src/dawn/native/Device.cpp +++ b/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 diff --git a/src/dawn/native/Device.h b/src/dawn/native/Device.h index 6bfb5ff8f5..de7e438cee 100644 --- a/src/dawn/native/Device.h +++ b/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; diff --git a/src/dawn/native/EncodingContext.cpp b/src/dawn/native/EncodingContext.cpp index b9ba529863..973cdecb50 100644 --- a/src/dawn/native/EncodingContext.cpp +++ b/src/dawn/native/EncodingContext.cpp @@ -91,12 +91,14 @@ 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. + if (mDevice->IsValidationEnabled() || + mDevice->MayRequireDuplicationOfIndirectParameters()) { + // When validation is enabled or indirect parameters require duplication, 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 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 diff --git a/src/dawn/native/IndirectDrawMetadata.cpp b/src/dawn/native/IndirectDrawMetadata.cpp index 3caf9f30a7..a1580efa3f 100644 --- a/src/dawn/native/IndirectDrawMetadata.cpp +++ b/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& + const std::vector& 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, - std::move(draw)); + it->second.AddIndirectDraw(mMaxDrawCallsPerBatch, mMaxBatchOffsetRange, 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 diff --git a/src/dawn/native/IndirectDrawMetadata.h b/src/dawn/native/IndirectDrawMetadata.h index 081c47b1fc..756a4ede91 100644 --- a/src/dawn/native/IndirectDrawMetadata.h +++ b/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 draws; + std::vector 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, - uint64_t maxBatchOffsetRange, - IndexedIndirectDraw draw); + void AddIndirectDraw(uint32_t maxDrawCallsPerIndirectValidationBatch, + uint64_t maxBatchOffsetRange, + 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& GetBatches() const; + const std::vector& GetBatches() const; private: Ref 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 mBatches; + std::vector 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; using IndexedIndirectBufferValidationInfoMap = std::map; @@ -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 mAddedBundles; diff --git a/src/dawn/native/IndirectDrawValidationEncoder.cpp b/src/dawn/native/IndirectDrawValidationEncoder.cpp index 7a9b14f940..98961ce4b0 100644 --- a/src/dawn/native/IndirectDrawValidationEncoder.cpp +++ b/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; - numIndexBufferElementsHigh: u32; - numDraws: u32; - padding: u32; - indirectOffsets: array; - }; + numIndexBufferElementsLow: u32, + numIndexBufferElementsHigh: u32, + numDraws: u32, + flags: u32, + indirectOffsets: array, + } struct IndirectParams { - data: array; - }; + data: array, + } @group(0) @binding(0) var batch: BatchInfo; @group(0) @binding(1) var clientParams: IndirectParams; @group(0) @binding(2) var 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; - 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; + let numParams = numIndirectParamsPerDrawCallValidated(); + let index = drawIndex * numParams; + for(var i = 0u; i < numParams; i = i + 1u) { + validatedParams.data[index + i] = 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]; - 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]; + + // The first 2 parameter is reserved for the duplicated first/baseVertex and firstInstance + + if (bool(batch.flags & kDuplicateBaseVertexInstance)) { + // first/baseVertex and firstInstance are always last two parameters + let dupIndex = cIndex + numClientParams - 2u; + validatedParams.data[vIndex] = clientParams.data[dupIndex]; + validatedParams.data[vIndex + 1u] = clientParams.data[dupIndex + 1u]; + + 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; - for (const IndirectDrawMetadata::IndexedIndirectValidationBatch& batch : + const uint64_t indirectDrawCommandSize = + 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; - newPass.clientIndirectBuffer = clientIndirectBuffer; + Pass newPass{}; + 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(batch.metadata->draws.size()); + batch.batchInfo->flags = pass.flags; uint32_t* indirectOffsets = reinterpret_cast(batch.batchInfo + 1); uint64_t validatedParamsOffset = batch.validatedParamsOffset; @@ -318,8 +382,11 @@ namespace dawn::native { draw.cmd->indirectBuffer = validatedParamsBuffer.GetBuffer(); draw.cmd->indirectOffset = validatedParamsOffset; - - validatedParamsOffset += kDrawIndexedIndirectSize; + 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 - // 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. + // Finally, we can now encode our validation and duplication passes. Each pass first does a + // two WriteBuffer to get batch and pass data over to the GPU, followed by a single compute + // 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(pass.batchData.get()), diff --git a/src/dawn/native/RenderEncoderBase.cpp b/src/dawn/native/RenderEncoderBase.cpp index e2607fb0f5..87cd32a7d0 100644 --- a/src/dawn/native/RenderEncoderBase.cpp +++ b/src/dawn/native/RenderEncoderBase.cpp @@ -172,9 +172,28 @@ namespace dawn::native { } DrawIndirectCmd* cmd = allocator->Allocate(Command::DrawIndirect); - cmd->indirectBuffer = indirectBuffer; - cmd->indirectOffset = indirectOffset; + 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(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 - // will be updated to point to it. + // 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; @@ -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; diff --git a/src/dawn/native/d3d12/CommandBufferD3D12.cpp b/src/dawn/native/d3d12/CommandBufferD3D12.cpp index 8f5fc8d86d..e3c23c81d4 100644 --- a/src/dawn/native/d3d12/CommandBufferD3D12.cpp +++ b/src/dawn/native/d3d12/CommandBufferD3D12.cpp @@ -149,19 +149,10 @@ namespace dawn::native::d3d12 { if (!firstOffsetInfo.usesVertexIndex && !firstOffsetInfo.usesInstanceIndex) { return; } - std::array offsets{}; - 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; - } + std::array offsets{firstVertex, firstInstance}; 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 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 signature = - ToBackend(GetDevice())->GetDrawIndexedIndirectSignature(); + lastPipeline->GetDrawIndexedIndirectCommandSignature(); commandList->ExecuteIndirect(signature.Get(), 1, buffer->GetD3D12Resource(), draw->indirectOffset, nullptr, 0); break; diff --git a/src/dawn/native/d3d12/DeviceD3D12.cpp b/src/dawn/native/d3d12/DeviceD3D12.cpp index 722eb9a3a0..4122480c40 100644 --- a/src/dawn/native/d3d12/DeviceD3D12.cpp +++ b/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 diff --git a/src/dawn/native/d3d12/DeviceD3D12.h b/src/dawn/native/d3d12/DeviceD3D12.h index 9bcd2b9091..84f0cd4ce5 100644 --- a/src/dawn/native/d3d12/DeviceD3D12.h +++ b/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: diff --git a/src/dawn/native/d3d12/PipelineLayoutD3D12.cpp b/src/dawn/native/d3d12/PipelineLayoutD3D12.cpp index 8eef9baeb4..8f1dbed99c 100644 --- a/src/dawn/native/d3d12/PipelineLayoutD3D12.cpp +++ b/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 diff --git a/src/dawn/native/d3d12/PipelineLayoutD3D12.h b/src/dawn/native/d3d12/PipelineLayoutD3D12.h index 99d3036e49..d9bb664253 100644 --- a/src/dawn/native/d3d12/PipelineLayoutD3D12.h +++ b/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 mRootSignature; ComPtr mDispatchIndirectCommandSignatureWithNumWorkgroups; + ComPtr mDrawIndirectCommandSignatureWithInstanceVertexOffsets; + ComPtr + mDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets; }; } // namespace dawn::native::d3d12 diff --git a/src/dawn/native/d3d12/RenderPipelineD3D12.cpp b/src/dawn/native/d3d12/RenderPipelineD3D12.cpp index a3441fed09..85129eaba9 100644 --- a/src/dawn/native/d3d12/RenderPipelineD3D12.cpp +++ b/src/dawn/native/d3d12/RenderPipelineD3D12.cpp @@ -463,6 +463,24 @@ namespace dawn::native::d3d12 { SetDebugName(ToBackend(GetDevice()), GetPipelineState(), "Dawn_RenderPipeline", GetLabel()); } + ComPtr RenderPipeline::GetDrawIndirectCommandSignature() { + if (mFirstOffsetInfo.usesVertexIndex || mFirstOffsetInfo.usesInstanceIndex) { + return ToBackend(GetLayout()) + ->GetDrawIndirectCommandSignatureWithInstanceVertexOffsets(); + } + + return ToBackend(GetDevice())->GetDrawIndirectSignature(); + } + + ComPtr RenderPipeline::GetDrawIndexedIndirectCommandSignature() { + if (mFirstOffsetInfo.usesVertexIndex || mFirstOffsetInfo.usesInstanceIndex) { + return ToBackend(GetLayout()) + ->GetDrawIndexedIndirectCommandSignatureWithInstanceVertexOffsets(); + } + + return ToBackend(GetDevice())->GetDrawIndexedIndirectSignature(); + } + D3D12_INPUT_LAYOUT_DESC RenderPipeline::ComputeInputLayout( std::array* inputElementDescriptors) { unsigned int count = 0; diff --git a/src/dawn/native/d3d12/RenderPipelineD3D12.h b/src/dawn/native/d3d12/RenderPipelineD3D12.h index 049520deeb..7d64bab579 100644 --- a/src/dawn/native/d3d12/RenderPipelineD3D12.h +++ b/src/dawn/native/d3d12/RenderPipelineD3D12.h @@ -43,6 +43,10 @@ namespace dawn::native::d3d12 { // Dawn API void SetLabelImpl() override; + ComPtr GetDrawIndirectCommandSignature(); + + ComPtr GetDrawIndexedIndirectCommandSignature(); + private: ~RenderPipeline() override; diff --git a/src/dawn/tests/end2end/BufferZeroInitTests.cpp b/src/dawn/tests/end2end/BufferZeroInitTests.cpp index 0fe2b57872..aefc48b605 100644 --- a/src/dawn/tests/end2end/BufferZeroInitTests.cpp +++ b/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. { diff --git a/src/dawn/tests/end2end/DrawIndirectTests.cpp b/src/dawn/tests/end2end/DrawIndirectTests.cpp index e8782813ac..84ab2b4634 100644 --- a/src/dawn/tests/end2end/DrawIndirectTests.cpp +++ b/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); diff --git a/src/dawn/tests/end2end/FirstIndexOffsetTests.cpp b/src/dawn/tests/end2end/FirstIndexOffsetTests.cpp index aa84387a9e..9dc9e1928b 100644 --- a/src/dawn/tests/end2end/FirstIndexOffsetTests.cpp +++ b/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 { @@ -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( 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( 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 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 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}. - if (IsD3D12() && (mode == DrawMode::NonIndexedIndirect || mode == DrawMode::IndexedIndirect)) { - expected = {0, 0}; + + // Per the specification, if validation is enabled and indirect-first-instance is not enabled, + // Draw[Indexed]Indirect with firstInstance > 0 will be a no-op. The buffer should still have + // 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(), diff --git a/src/tint/transform/first_index_offset.cc b/src/tint/transform/first_index_offset.cc index 46a2dcd245..2419bc255d 100644 --- a/src/tint/transform/first_index_offset.cc +++ b/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; - offset += 4; - } - if (has_instance_index) { - members.push_back(ctx.dst->Member(kFirstInstanceName, ctx.dst->ty.u32())); - instance_index_offset = offset; - offset += 4; - } + members.push_back(ctx.dst->Member(kFirstVertexName, ctx.dst->ty.u32())); + vertex_index_offset = 0; + members.push_back(ctx.dst->Member(kFirstInstanceName, ctx.dst->ty.u32())); + instance_index_offset = 4; auto* struct_ = ctx.dst->Structure(ctx.dst->Sym(), std::move(members)); // Create a global to hold the uniform buffer diff --git a/webgpu-cts/expectations.txt b/webgpu-cts/expectations.txt index 6ae1323dd2..7e710dacd6 100644 --- a/webgpu-cts/expectations.txt +++ b/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