diff --git a/src/dawn/native/CommandEncoder.cpp b/src/dawn/native/CommandEncoder.cpp index 0bc63fa07d..db36bd66af 100644 --- a/src/dawn/native/CommandEncoder.cpp +++ b/src/dawn/native/CommandEncoder.cpp @@ -778,8 +778,6 @@ ComputePassEncoder* CommandEncoder::APIBeginComputePass(const ComputePassDescrip Ref CommandEncoder::BeginComputePass(const ComputePassDescriptor* descriptor) { DeviceBase* device = GetDevice(); - std::vector timestampWritesAtBeginning; - std::vector timestampWritesAtEnd; bool success = mEncodingContext.TryEncode( this, [&](CommandAllocator* allocator) -> MaybeError { @@ -792,27 +790,27 @@ Ref CommandEncoder::BeginComputePass(const ComputePassDescri return {}; } - // Split the timestampWrites used in BeginComputePassCmd and EndComputePassCmd + // Record timestamp writes at the beginning and end of compute pass. The timestamp write + // at the end also be needed in BeginComputePassCmd because it's required by compute + // pass descriptor when beginning compute pass on Metal. for (uint32_t i = 0; i < descriptor->timestampWriteCount; i++) { QuerySetBase* querySet = descriptor->timestampWrites[i].querySet; uint32_t queryIndex = descriptor->timestampWrites[i].queryIndex; switch (descriptor->timestampWrites[i].location) { case wgpu::ComputePassTimestampLocation::Beginning: - timestampWritesAtBeginning.push_back({querySet, queryIndex}); + cmd->beginTimestamp.querySet = querySet; + cmd->beginTimestamp.queryIndex = queryIndex; break; case wgpu::ComputePassTimestampLocation::End: - timestampWritesAtEnd.push_back({querySet, queryIndex}); - break; - default: + cmd->endTimestamp.querySet = querySet; + cmd->endTimestamp.queryIndex = queryIndex; break; } TrackQueryAvailability(querySet, queryIndex); } - cmd->timestampWrites = std::move(timestampWritesAtBeginning); - return {}; }, "encoding %s.BeginComputePass(%s).", this, descriptor); @@ -823,8 +821,8 @@ Ref CommandEncoder::BeginComputePass(const ComputePassDescri descriptor = &defaultDescriptor; } - Ref passEncoder = ComputePassEncoder::Create( - device, descriptor, this, &mEncodingContext, std::move(timestampWritesAtEnd)); + Ref passEncoder = + ComputePassEncoder::Create(device, descriptor, this, &mEncodingContext); mEncodingContext.EnterPass(passEncoder.Get()); return passEncoder; } @@ -846,8 +844,6 @@ Ref CommandEncoder::BeginRenderPass(const RenderPassDescripto bool depthReadOnly = false; bool stencilReadOnly = false; Ref attachmentState; - std::vector timestampWritesAtBeginning; - std::vector timestampWritesAtEnd; bool success = mEncodingContext.TryEncode( this, [&](CommandAllocator* allocator) -> MaybeError { @@ -865,28 +861,6 @@ Ref CommandEncoder::BeginRenderPass(const RenderPassDescripto cmd->attachmentState = device->GetOrCreateAttachmentState(descriptor); attachmentState = cmd->attachmentState; - // Split the timestampWrites used in BeginRenderPassCmd and EndRenderPassCmd - for (uint32_t i = 0; i < descriptor->timestampWriteCount; i++) { - QuerySetBase* querySet = descriptor->timestampWrites[i].querySet; - uint32_t queryIndex = descriptor->timestampWrites[i].queryIndex; - - switch (descriptor->timestampWrites[i].location) { - case wgpu::RenderPassTimestampLocation::Beginning: - timestampWritesAtBeginning.push_back({querySet, queryIndex}); - break; - case wgpu::RenderPassTimestampLocation::End: - timestampWritesAtEnd.push_back({querySet, queryIndex}); - break; - default: - break; - } - - TrackQueryAvailability(querySet, queryIndex); - // Track the query availability with true on render pass again for rewrite - // validation and query reset on Vulkan - usageTracker.TrackQueryAvailability(querySet, queryIndex); - } - for (ColorAttachmentIndex index : IterateBitSet(cmd->attachmentState->GetColorAttachmentsMask())) { uint8_t i = static_cast(index); @@ -990,7 +964,29 @@ Ref CommandEncoder::BeginRenderPass(const RenderPassDescripto cmd->occlusionQuerySet = descriptor->occlusionQuerySet; - cmd->timestampWrites = std::move(timestampWritesAtBeginning); + // Record timestamp writes at the beginning and end of render pass. The timestamp write + // at the end also be needed in BeginComputePassCmd because it's required by render pass + // descriptor when beginning render pass on Metal. + for (uint32_t i = 0; i < descriptor->timestampWriteCount; i++) { + QuerySetBase* querySet = descriptor->timestampWrites[i].querySet; + uint32_t queryIndex = descriptor->timestampWrites[i].queryIndex; + + switch (descriptor->timestampWrites[i].location) { + case wgpu::RenderPassTimestampLocation::Beginning: + cmd->beginTimestamp.querySet = querySet; + cmd->beginTimestamp.queryIndex = queryIndex; + break; + case wgpu::RenderPassTimestampLocation::End: + cmd->endTimestamp.querySet = querySet; + cmd->endTimestamp.queryIndex = queryIndex; + break; + } + + TrackQueryAvailability(querySet, queryIndex); + // Track the query availability with true on render pass again for rewrite + // validation and query reset on Vulkan + usageTracker.TrackQueryAvailability(querySet, queryIndex); + } return {}; }, @@ -999,8 +995,7 @@ Ref CommandEncoder::BeginRenderPass(const RenderPassDescripto if (success) { Ref passEncoder = RenderPassEncoder::Create( device, descriptor, this, &mEncodingContext, std::move(usageTracker), - std::move(attachmentState), std::move(timestampWritesAtEnd), width, height, - depthReadOnly, stencilReadOnly); + std::move(attachmentState), width, height, depthReadOnly, stencilReadOnly); mEncodingContext.EnterPass(passEncoder.Get()); return passEncoder; } diff --git a/src/dawn/native/Commands.cpp b/src/dawn/native/Commands.cpp index 6f3d6df467..c1be63922f 100644 --- a/src/dawn/native/Commands.cpp +++ b/src/dawn/native/Commands.cpp @@ -361,9 +361,7 @@ void SkipCommand(CommandIterator* commands, Command type) { } } -TimestampWrite::TimestampWrite(const Ref& set, uint32_t idx) - : querySet(set), queryIndex(idx) {} -TimestampWrite::TimestampWrite(TimestampWrite&&) = default; +TimestampWrite::TimestampWrite() = default; TimestampWrite::~TimestampWrite() = default; BeginComputePassCmd::BeginComputePassCmd() = default; diff --git a/src/dawn/native/Commands.h b/src/dawn/native/Commands.h index c7bfa04211..456a1b5961 100644 --- a/src/dawn/native/Commands.h +++ b/src/dawn/native/Commands.h @@ -70,8 +70,7 @@ enum class Command { }; struct TimestampWrite { - TimestampWrite(const Ref& set, uint32_t idx); - TimestampWrite(TimestampWrite&&); + TimestampWrite(); ~TimestampWrite(); Ref querySet; @@ -82,7 +81,8 @@ struct BeginComputePassCmd { BeginComputePassCmd(); ~BeginComputePassCmd(); - std::vector timestampWrites; + TimestampWrite beginTimestamp; + TimestampWrite endTimestamp; }; struct BeginOcclusionQueryCmd { @@ -133,7 +133,8 @@ struct BeginRenderPassCmd { uint32_t height; Ref occlusionQuerySet; - std::vector timestampWrites; + TimestampWrite beginTimestamp; + TimestampWrite endTimestamp; }; struct BufferCopy { @@ -228,8 +229,6 @@ struct DrawIndexedIndirectCmd : DrawIndirectCmd {}; struct EndComputePassCmd { EndComputePassCmd(); ~EndComputePassCmd(); - - std::vector timestampWrites; }; struct EndOcclusionQueryCmd { @@ -243,8 +242,6 @@ struct EndOcclusionQueryCmd { struct EndRenderPassCmd { EndRenderPassCmd(); ~EndRenderPassCmd(); - - std::vector timestampWrites; }; struct ExecuteBundlesCmd { diff --git a/src/dawn/native/ComputePassEncoder.cpp b/src/dawn/native/ComputePassEncoder.cpp index 30fa36bcab..65d00cbc91 100644 --- a/src/dawn/native/ComputePassEncoder.cpp +++ b/src/dawn/native/ComputePassEncoder.cpp @@ -111,23 +111,18 @@ ResultOrError GetOrCreateIndirectDispatchValidationPipelin ComputePassEncoder::ComputePassEncoder(DeviceBase* device, const ComputePassDescriptor* descriptor, CommandEncoder* commandEncoder, - EncodingContext* encodingContext, - std::vector timestampWritesAtEnd) + EncodingContext* encodingContext) : ProgrammableEncoder(device, descriptor->label, encodingContext), - mCommandEncoder(commandEncoder), - mTimestampWritesAtEnd(std::move(timestampWritesAtEnd)) { + mCommandEncoder(commandEncoder) { TrackInDevice(); } // static -Ref ComputePassEncoder::Create( - DeviceBase* device, - const ComputePassDescriptor* descriptor, - CommandEncoder* commandEncoder, - EncodingContext* encodingContext, - std::vector timestampWritesAtEnd) { - return AcquireRef(new ComputePassEncoder(device, descriptor, commandEncoder, encodingContext, - std::move(timestampWritesAtEnd))); +Ref ComputePassEncoder::Create(DeviceBase* device, + const ComputePassDescriptor* descriptor, + CommandEncoder* commandEncoder, + EncodingContext* encodingContext) { + return AcquireRef(new ComputePassEncoder(device, descriptor, commandEncoder, encodingContext)); } ComputePassEncoder::ComputePassEncoder(DeviceBase* device, @@ -162,11 +157,7 @@ void ComputePassEncoder::APIEnd() { DAWN_TRY(ValidateProgrammableEncoderEnd()); } - EndComputePassCmd* cmd = - allocator->Allocate(Command::EndComputePass); - // The query availability has already been updated at the beginning of compute - // pass, and no need to do update here. - cmd->timestampWrites = std::move(mTimestampWritesAtEnd); + allocator->Allocate(Command::EndComputePass); return {}; }, diff --git a/src/dawn/native/ComputePassEncoder.h b/src/dawn/native/ComputePassEncoder.h index ad950964b2..f48ef94cb7 100644 --- a/src/dawn/native/ComputePassEncoder.h +++ b/src/dawn/native/ComputePassEncoder.h @@ -33,8 +33,7 @@ class ComputePassEncoder final : public ProgrammableEncoder { static Ref Create(DeviceBase* device, const ComputePassDescriptor* descriptor, CommandEncoder* commandEncoder, - EncodingContext* encodingContext, - std::vector timestampWritesAtEnd); + EncodingContext* encodingContext); static Ref MakeError(DeviceBase* device, CommandEncoder* commandEncoder, EncodingContext* encodingContext); @@ -72,8 +71,7 @@ class ComputePassEncoder final : public ProgrammableEncoder { ComputePassEncoder(DeviceBase* device, const ComputePassDescriptor* descriptor, CommandEncoder* commandEncoder, - EncodingContext* encodingContext, - std::vector timestampWritesAtEnd); + EncodingContext* encodingContext); ComputePassEncoder(DeviceBase* device, CommandEncoder* commandEncoder, EncodingContext* encodingContext, @@ -98,8 +96,6 @@ class ComputePassEncoder final : public ProgrammableEncoder { // For render and compute passes, the encoding context is borrowed from the command encoder. // Keep a reference to the encoder to make sure the context isn't freed. Ref mCommandEncoder; - - std::vector mTimestampWritesAtEnd; }; } // namespace dawn::native diff --git a/src/dawn/native/RenderPassEncoder.cpp b/src/dawn/native/RenderPassEncoder.cpp index 9994cf5941..b066655f37 100644 --- a/src/dawn/native/RenderPassEncoder.cpp +++ b/src/dawn/native/RenderPassEncoder.cpp @@ -56,7 +56,6 @@ RenderPassEncoder::RenderPassEncoder(DeviceBase* device, EncodingContext* encodingContext, RenderPassResourceUsageTracker usageTracker, Ref attachmentState, - std::vector timestampWritesAtEnd, uint32_t renderTargetWidth, uint32_t renderTargetHeight, bool depthReadOnly, @@ -70,8 +69,7 @@ RenderPassEncoder::RenderPassEncoder(DeviceBase* device, mCommandEncoder(commandEncoder), mRenderTargetWidth(renderTargetWidth), mRenderTargetHeight(renderTargetHeight), - mOcclusionQuerySet(descriptor->occlusionQuerySet), - mTimestampWritesAtEnd(std::move(timestampWritesAtEnd)) { + mOcclusionQuerySet(descriptor->occlusionQuerySet) { mUsageTracker = std::move(usageTracker); const RenderPassDescriptorMaxDrawCount* maxDrawCountInfo = nullptr; FindInChain(descriptor->nextInChain, &maxDrawCountInfo); @@ -88,15 +86,14 @@ Ref RenderPassEncoder::Create(DeviceBase* device, EncodingContext* encodingContext, RenderPassResourceUsageTracker usageTracker, Ref attachmentState, - std::vector timestampWritesAtEnd, uint32_t renderTargetWidth, uint32_t renderTargetHeight, bool depthReadOnly, bool stencilReadOnly) { return AcquireRef(new RenderPassEncoder(device, descriptor, commandEncoder, encodingContext, std::move(usageTracker), std::move(attachmentState), - std::move(timestampWritesAtEnd), renderTargetWidth, - renderTargetHeight, depthReadOnly, stencilReadOnly)); + renderTargetWidth, renderTargetHeight, depthReadOnly, + stencilReadOnly)); } RenderPassEncoder::RenderPassEncoder(DeviceBase* device, @@ -152,10 +149,7 @@ void RenderPassEncoder::APIEnd() { mDrawCount, this, mMaxDrawCount); } - EndRenderPassCmd* cmd = allocator->Allocate(Command::EndRenderPass); - // The query availability has already been updated at the beginning of render - // pass, and no need to do update here. - cmd->timestampWrites = std::move(mTimestampWritesAtEnd); + allocator->Allocate(Command::EndRenderPass); DAWN_TRY(mEncodingContext->ExitRenderPass(this, std::move(mUsageTracker), mCommandEncoder.Get(), diff --git a/src/dawn/native/RenderPassEncoder.h b/src/dawn/native/RenderPassEncoder.h index 45714d0849..32199f1654 100644 --- a/src/dawn/native/RenderPassEncoder.h +++ b/src/dawn/native/RenderPassEncoder.h @@ -33,7 +33,6 @@ class RenderPassEncoder final : public RenderEncoderBase { EncodingContext* encodingContext, RenderPassResourceUsageTracker usageTracker, Ref attachmentState, - std::vector timestampWritesAtEnd, uint32_t renderTargetWidth, uint32_t renderTargetHeight, bool depthReadOnly, @@ -70,7 +69,6 @@ class RenderPassEncoder final : public RenderEncoderBase { EncodingContext* encodingContext, RenderPassResourceUsageTracker usageTracker, Ref attachmentState, - std::vector timestampWritesAtEnd, uint32_t renderTargetWidth, uint32_t renderTargetHeight, bool depthReadOnly, @@ -99,8 +97,6 @@ class RenderPassEncoder final : public RenderEncoderBase { // This is the hardcoded value in the WebGPU spec. uint64_t mMaxDrawCount = 50000000; - - std::vector mTimestampWritesAtEnd; }; } // namespace dawn::native diff --git a/src/dawn/native/d3d12/CommandBufferD3D12.cpp b/src/dawn/native/d3d12/CommandBufferD3D12.cpp index 77c7ba3ca7..08e2bcbe16 100644 --- a/src/dawn/native/d3d12/CommandBufferD3D12.cpp +++ b/src/dawn/native/d3d12/CommandBufferD3D12.cpp @@ -90,10 +90,12 @@ bool CanUseCopyResource(const TextureCopy& src, const TextureCopy& dst, const Ex copySize.depthOrArrayLayers == srcSize.depthOrArrayLayers; } -void RecordWriteTimestampCmd(ID3D12GraphicsCommandList* commandList, WriteTimestampCmd* cmd) { - QuerySet* querySet = ToBackend(cmd->querySet.Get()); - ASSERT(D3D12QueryType(querySet->GetQueryType()) == D3D12_QUERY_TYPE_TIMESTAMP); - commandList->EndQuery(querySet->GetQueryHeap(), D3D12_QUERY_TYPE_TIMESTAMP, cmd->queryIndex); +void RecordWriteTimestampCmd(ID3D12GraphicsCommandList* commandList, + QuerySetBase* querySet, + uint32_t queryIndex) { + ASSERT(D3D12QueryType(ToBackend(querySet)->GetQueryType()) == D3D12_QUERY_TYPE_TIMESTAMP); + commandList->EndQuery(ToBackend(querySet)->GetQueryHeap(), D3D12_QUERY_TYPE_TIMESTAMP, + queryIndex); } void RecordResolveQuerySetCmd(ID3D12GraphicsCommandList* commandList, @@ -653,11 +655,12 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* commandContext while (mCommands.NextCommandId(&type)) { switch (type) { case Command::BeginComputePass: { - mCommands.NextCommand(); + BeginComputePassCmd* cmd = mCommands.NextCommand(); bindingTracker.SetInComputePass(true); + DAWN_TRY( - RecordComputePass(commandContext, &bindingTracker, + RecordComputePass(commandContext, &bindingTracker, cmd, GetResourceUsages().computePasses[nextComputePassNumber])); nextComputePassNumber++; @@ -942,7 +945,7 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* commandContext case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - RecordWriteTimestampCmd(commandList, cmd); + RecordWriteTimestampCmd(commandList, cmd->querySet.Get(), cmd->queryIndex); break; } @@ -1023,10 +1026,17 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* commandContext MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* commandContext, BindGroupStateTracker* bindingTracker, + BeginComputePassCmd* computePass, const ComputePassResourceUsage& resourceUsages) { uint64_t currentDispatch = 0; ID3D12GraphicsCommandList* commandList = commandContext->GetCommandList(); + // Write timestamp at the beginning of compute pass if it's set. + if (computePass->beginTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(commandList, computePass->beginTimestamp.querySet.Get(), + computePass->beginTimestamp.queryIndex); + } + Command type; ComputePipeline* lastPipeline = nullptr; while (mCommands.NextCommandId(&type)) { @@ -1068,6 +1078,12 @@ MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* commandCont case Command::EndComputePass: { mCommands.NextCommand(); + + // Write timestamp at the end of compute pass if it's set. + if (computePass->endTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(commandList, computePass->endTimestamp.querySet.Get(), + computePass->endTimestamp.queryIndex); + } return {}; } @@ -1136,7 +1152,7 @@ MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* commandCont case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - RecordWriteTimestampCmd(commandList, cmd); + RecordWriteTimestampCmd(commandList, cmd->querySet.Get(), cmd->queryIndex); break; } @@ -1339,6 +1355,12 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* commandConte ID3D12GraphicsCommandList* commandList = commandContext->GetCommandList(); + // Write timestamp at the beginning of render pass if it's set. + if (renderPass->beginTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(commandList, renderPass->beginTimestamp.querySet.Get(), + renderPass->beginTimestamp.queryIndex); + } + // Set up default dynamic state { uint32_t width = renderPass->width; @@ -1511,6 +1533,13 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* commandConte switch (type) { case Command::EndRenderPass: { mCommands.NextCommand(); + + // Write timestamp at the end of render pass if it's set. + if (renderPass->endTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(commandList, renderPass->endTimestamp.querySet.Get(), + renderPass->endTimestamp.queryIndex); + } + if (useRenderPass) { commandContext->GetCommandList4()->EndRenderPass(); } else if (renderPass->attachmentState->GetSampleCount() > 1) { @@ -1596,7 +1625,7 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* commandConte case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - RecordWriteTimestampCmd(commandList, cmd); + RecordWriteTimestampCmd(commandList, cmd->querySet.Get(), cmd->queryIndex); break; } diff --git a/src/dawn/native/d3d12/CommandBufferD3D12.h b/src/dawn/native/d3d12/CommandBufferD3D12.h index ea0bb704ac..10523950a8 100644 --- a/src/dawn/native/d3d12/CommandBufferD3D12.h +++ b/src/dawn/native/d3d12/CommandBufferD3D12.h @@ -19,6 +19,7 @@ #include "dawn/native/Error.h" namespace dawn::native { +struct BeginComputePassCmd; struct BeginRenderPassCmd; } // namespace dawn::native @@ -40,6 +41,7 @@ class CommandBuffer final : public CommandBufferBase { MaybeError RecordComputePass(CommandRecordingContext* commandContext, BindGroupStateTracker* bindingTracker, + BeginComputePassCmd* computePass, const ComputePassResourceUsage& resourceUsages); MaybeError RecordRenderPass(CommandRecordingContext* commandContext, BindGroupStateTracker* bindingTracker, diff --git a/src/dawn/native/metal/CommandBufferMTL.mm b/src/dawn/native/metal/CommandBufferMTL.mm index d32f058911..66a4569277 100644 --- a/src/dawn/native/metal/CommandBufferMTL.mm +++ b/src/dawn/native/metal/CommandBufferMTL.mm @@ -610,7 +610,12 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext) while (mCommands.NextCommandId(&type)) { switch (type) { case Command::BeginComputePass: { - mCommands.NextCommand(); + BeginComputePassCmd* cmd = mCommands.NextCommand(); + + if (cmd->beginTimestamp.querySet.Get() != nullptr || + cmd->endTimestamp.querySet.Get() != nullptr) { + return DAWN_UNIMPLEMENTED_ERROR("timestampWrites unimplemented."); + } for (const SyncScopeResourceUsage& scope : GetResourceUsages().computePasses[nextComputePassNumber].dispatchUsages) { @@ -627,6 +632,11 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext) case Command::BeginRenderPass: { BeginRenderPassCmd* cmd = mCommands.NextCommand(); + if (cmd->beginTimestamp.querySet.Get() != nullptr || + cmd->endTimestamp.querySet.Get() != nullptr) { + return DAWN_UNIMPLEMENTED_ERROR("timestampWrites unimplemented."); + } + LazyClearSyncScope(GetResourceUsages().renderPasses[nextRenderPassNumber], commandContext); commandContext->EndBlit(); diff --git a/src/dawn/native/vulkan/CommandBufferVk.cpp b/src/dawn/native/vulkan/CommandBufferVk.cpp index fca86c2db7..a61a800d98 100644 --- a/src/dawn/native/vulkan/CommandBufferVk.cpp +++ b/src/dawn/native/vulkan/CommandBufferVk.cpp @@ -377,12 +377,19 @@ void ResetUsedQuerySetsOnRenderPass(Device* device, void RecordWriteTimestampCmd(CommandRecordingContext* recordingContext, Device* device, - WriteTimestampCmd* cmd) { + QuerySetBase* querySet, + uint32_t queryIndex, + bool isRenderPass) { VkCommandBuffer commands = recordingContext->commandBuffer; - QuerySet* querySet = ToBackend(cmd->querySet.Get()); + + // The queries must be reset between uses, and the reset command cannot be called in render + // pass. + if (!isRenderPass) { + device->fn.CmdResetQueryPool(commands, ToBackend(querySet)->GetHandle(), queryIndex, 1); + } device->fn.CmdWriteTimestamp(commands, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, - querySet->GetHandle(), cmd->queryIndex); + ToBackend(querySet)->GetHandle(), queryIndex); } void RecordResolveQuerySetCmd(VkCommandBuffer commands, @@ -735,10 +742,11 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* recordingConte } case Command::BeginComputePass: { - mCommands.NextCommand(); + BeginComputePassCmd* cmd = mCommands.NextCommand(); - DAWN_TRY(RecordComputePass( - recordingContext, GetResourceUsages().computePasses[nextComputePassNumber])); + DAWN_TRY( + RecordComputePass(recordingContext, cmd, + GetResourceUsages().computePasses[nextComputePassNumber])); nextComputePassNumber++; break; @@ -777,11 +785,8 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* recordingConte case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - // The query must be reset between uses. - device->fn.CmdResetQueryPool(commands, ToBackend(cmd->querySet)->GetHandle(), - cmd->queryIndex, 1); - - RecordWriteTimestampCmd(recordingContext, device, cmd); + RecordWriteTimestampCmd(recordingContext, device, cmd->querySet.Get(), + cmd->queryIndex, false); break; } @@ -878,8 +883,17 @@ MaybeError CommandBuffer::RecordCommands(CommandRecordingContext* recordingConte } MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* recordingContext, + BeginComputePassCmd* computePassCmd, const ComputePassResourceUsage& resourceUsages) { Device* device = ToBackend(GetDevice()); + + // Write timestamp at the beginning of compute pass if it's set + if (computePassCmd->beginTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(recordingContext, device, + computePassCmd->beginTimestamp.querySet.Get(), + computePassCmd->beginTimestamp.queryIndex, false); + } + VkCommandBuffer commands = recordingContext->commandBuffer; uint64_t currentDispatch = 0; @@ -890,6 +904,13 @@ MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* recordingCo switch (type) { case Command::EndComputePass: { mCommands.NextCommand(); + + // Write timestamp at the end of compute pass if it's set. + if (computePassCmd->endTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(recordingContext, device, + computePassCmd->endTimestamp.querySet.Get(), + computePassCmd->endTimestamp.queryIndex, false); + } return {}; } @@ -996,11 +1017,8 @@ MaybeError CommandBuffer::RecordComputePass(CommandRecordingContext* recordingCo case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - // The query must be reset between uses. - device->fn.CmdResetQueryPool(commands, ToBackend(cmd->querySet)->GetHandle(), - cmd->queryIndex, 1); - - RecordWriteTimestampCmd(recordingContext, device, cmd); + RecordWriteTimestampCmd(recordingContext, device, cmd->querySet.Get(), + cmd->queryIndex, false); break; } @@ -1020,6 +1038,13 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* recordingCon DAWN_TRY(RecordBeginRenderPass(recordingContext, device, renderPassCmd)); + // Write timestamp at the beginning of render pass if it's set. + if (renderPassCmd->beginTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(recordingContext, device, + renderPassCmd->beginTimestamp.querySet.Get(), + renderPassCmd->beginTimestamp.queryIndex, true); + } + // Set the default value for the dynamic state { device->fn.CmdSetLineWidth(commands, 1.0f); @@ -1203,6 +1228,14 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* recordingCon switch (type) { case Command::EndRenderPass: { mCommands.NextCommand(); + + // Write timestamp at the end of render pass if it's set. + if (renderPassCmd->endTimestamp.querySet.Get() != nullptr) { + RecordWriteTimestampCmd(recordingContext, device, + renderPassCmd->endTimestamp.querySet.Get(), + renderPassCmd->endTimestamp.queryIndex, true); + } + device->fn.CmdEndRenderPass(commands); return {}; } @@ -1290,7 +1323,8 @@ MaybeError CommandBuffer::RecordRenderPass(CommandRecordingContext* recordingCon case Command::WriteTimestamp: { WriteTimestampCmd* cmd = mCommands.NextCommand(); - RecordWriteTimestampCmd(recordingContext, device, cmd); + RecordWriteTimestampCmd(recordingContext, device, cmd->querySet.Get(), + cmd->queryIndex, true); break; } diff --git a/src/dawn/native/vulkan/CommandBufferVk.h b/src/dawn/native/vulkan/CommandBufferVk.h index dbb7fdcd7f..b748f0daac 100644 --- a/src/dawn/native/vulkan/CommandBufferVk.h +++ b/src/dawn/native/vulkan/CommandBufferVk.h @@ -21,6 +21,7 @@ #include "dawn/common/vulkan_platform.h" namespace dawn::native { +struct BeginComputePassCmd; struct BeginRenderPassCmd; struct TextureCopy; } // namespace dawn::native @@ -41,6 +42,7 @@ class CommandBuffer final : public CommandBufferBase { CommandBuffer(CommandEncoder* encoder, const CommandBufferDescriptor* descriptor); MaybeError RecordComputePass(CommandRecordingContext* recordingContext, + BeginComputePassCmd* computePass, const ComputePassResourceUsage& resourceUsages); MaybeError RecordRenderPass(CommandRecordingContext* recordingContext, BeginRenderPassCmd* renderPass); diff --git a/src/dawn/tests/end2end/QueryTests.cpp b/src/dawn/tests/end2end/QueryTests.cpp index 90e3467483..ff195376c6 100644 --- a/src/dawn/tests/end2end/QueryTests.cpp +++ b/src/dawn/tests/end2end/QueryTests.cpp @@ -302,7 +302,7 @@ TEST_P(OcclusionQueryTests, Rewrite) { // the query resetting at the start of render passes on Vulkan backend. TEST_P(OcclusionQueryTests, ResolveSparseQueries) { // TODO(hao.x.li@intel.com): Fails on Intel Windows Vulkan due to a driver issue that - // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it util + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until // the issue is fixed. DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel()); @@ -523,6 +523,17 @@ class TimestampQueryTests : public QueryTests { // Skip all tests if timestamp feature is not supported DAWN_TEST_UNSUPPORTED_IF(!SupportsFeatures({wgpu::FeatureName::TimestampQuery})); + + // Create basic compute pipeline + wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( + @compute @workgroup_size(1) + fn main() { + })"); + + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = module; + csDesc.compute.entryPoint = "main"; + computePipeline = device.CreateComputePipeline(&csDesc); } std::vector GetRequiredFeatures() override { @@ -539,6 +550,113 @@ class TimestampQueryTests : public QueryTests { descriptor.type = wgpu::QueryType::Timestamp; return device.CreateQuerySet(&descriptor); } + + void TestTimestampWritesOnComputePass( + const std::vector& timestampWrites, + const std::vector& timestampWritesOnAnotherPass = {}) { + size_t queryCount = timestampWrites.size() + timestampWritesOnAnotherPass.size(); + // The destination buffer offset must be a multiple of 256. + wgpu::Buffer destination = + CreateResolveBuffer(queryCount * kMinDestinationOffset + sizeof(uint64_t)); + + wgpu::ComputePassDescriptor descriptor; + descriptor.timestampWriteCount = timestampWrites.size(); + descriptor.timestampWrites = timestampWrites.data(); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(&descriptor); + pass.SetPipeline(computePipeline); + pass.DispatchWorkgroups(1, 1, 1); + pass.End(); + + // Resolve queries one by one because the query set at the beginning of pass may be + // different with the one at the end of pass. + for (size_t i = 0; i < timestampWrites.size(); i++) { + encoder.ResolveQuerySet(timestampWrites[i].querySet, timestampWrites[i].queryIndex, 1, + destination, i * kMinDestinationOffset); + } + + // Begin another compute pass if the timestampWritesOnAnotherPass is set. + if (!timestampWritesOnAnotherPass.empty()) { + wgpu::ComputePassDescriptor descriptor2; + descriptor2.timestampWriteCount = timestampWritesOnAnotherPass.size(); + descriptor2.timestampWrites = timestampWritesOnAnotherPass.data(); + + wgpu::ComputePassEncoder pass2 = encoder.BeginComputePass(&descriptor2); + pass2.SetPipeline(computePipeline); + pass2.DispatchWorkgroups(1, 1, 1); + pass2.End(); + + for (size_t i = 0; i < timestampWritesOnAnotherPass.size(); i++) { + // Resolve queries one by one because the query set at the beginning of pass may be + // different with the one at the end of pass. + encoder.ResolveQuerySet(timestampWritesOnAnotherPass[i].querySet, + timestampWritesOnAnotherPass[i].queryIndex, 1, destination, + (timestampWrites.size() + i) * kMinDestinationOffset); + } + } + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + for (size_t i = 0; i < queryCount; i++) { + EXPECT_BUFFER(destination, i * kMinDestinationOffset, sizeof(uint64_t), + new TimestampExpectation); + } + } + + void TestTimestampWritesOnRenderPass( + const std::vector& timestampWrites, + const std::vector& timestampWritesOnAnotherPass = {}) { + size_t queryCount = timestampWrites.size() + timestampWritesOnAnotherPass.size(); + // The destination buffer offset must be a multiple of 256. + wgpu::Buffer destination = + CreateResolveBuffer(queryCount * kMinDestinationOffset + sizeof(uint64_t)); + + utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); + renderPass.renderPassInfo.timestampWriteCount = timestampWrites.size(); + renderPass.renderPassInfo.timestampWrites = timestampWrites.data(); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.End(); + + // Resolve queries one by one because the query set at the beginning of pass may be + // different with the one at the end of pass. + for (size_t i = 0; i < timestampWrites.size(); i++) { + encoder.ResolveQuerySet(timestampWrites[i].querySet, timestampWrites[i].queryIndex, 1, + destination, i * kMinDestinationOffset); + } + + // Begin another render pass if the timestampWritesOnAnotherPass is set. + if (!timestampWritesOnAnotherPass.empty()) { + utils::BasicRenderPass renderPass2 = utils::CreateBasicRenderPass(device, 1, 1); + renderPass2.renderPassInfo.timestampWriteCount = timestampWritesOnAnotherPass.size(); + renderPass2.renderPassInfo.timestampWrites = timestampWritesOnAnotherPass.data(); + + wgpu::RenderPassEncoder pass2 = encoder.BeginRenderPass(&renderPass2.renderPassInfo); + pass2.End(); + + for (size_t i = 0; i < timestampWritesOnAnotherPass.size(); i++) { + // Resolve queries one by one because the query set at the beginning of pass may be + // different with the one at the end of pass. + encoder.ResolveQuerySet(timestampWritesOnAnotherPass[i].querySet, + timestampWritesOnAnotherPass[i].queryIndex, 1, destination, + (timestampWrites.size() + i) * kMinDestinationOffset); + } + } + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + for (size_t i = 0; i < queryCount; i++) { + EXPECT_BUFFER(destination, i * kMinDestinationOffset, sizeof(uint64_t), + new TimestampExpectation); + } + } + + private: + wgpu::ComputePipeline computePipeline; }; // Test creating query set with the type of Timestamp @@ -551,6 +669,9 @@ TEST_P(TimestampQueryTests, QuerySetCreation) { // Test calling timestamp query from command encoder TEST_P(TimestampQueryTests, TimestampOnCommandEncoder) { + // TODO (dawn:1250): Still not implemented on Metal backend. + DAWN_TEST_UNSUPPORTED_IF(IsMetal()); + constexpr uint32_t kQueryCount = 2; // Write timestamp with different query indexes @@ -694,6 +815,124 @@ TEST_P(TimestampQueryTests, TimestampOnComputePass) { } } +// Test timestampWrites setting in compute pass descriptor +TEST_P(TimestampQueryTests, TimestampWritesOnComputePass) { + // TODO(dawn:1489): Fails on Intel Windows Vulkan due to a driver issue that + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until + // the issue is fixed. + DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel()); + + // TODO (dawn:1250): Still not implemented on Metal backend. + DAWN_TEST_UNSUPPORTED_IF(IsMetal()); + + constexpr uint32_t kQueryCount = 2; + + // Set timestampWrites with different query indexes and locations on same compute pass + { + wgpu::QuerySet querySet = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnComputePass( + {{querySet, 0, wgpu::ComputePassTimestampLocation::Beginning}, + {querySet, 1, wgpu::ComputePassTimestampLocation::End}}); + } + + // Set timestampWrites with different query set on same compute pass + { + wgpu::QuerySet querySet0 = CreateQuerySetForTimestamp(1); + wgpu::QuerySet querySet1 = CreateQuerySetForTimestamp(1); + + TestTimestampWritesOnComputePass( + {{querySet0, 0, wgpu::ComputePassTimestampLocation::Beginning}, + {querySet1, 0, wgpu::ComputePassTimestampLocation::End}}); + } + + // Set timestampWrites with only one value of ComputePassTimestampLocation + { + wgpu::QuerySet querySet = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnComputePass( + {{querySet, 0, wgpu::ComputePassTimestampLocation::Beginning}}); + + TestTimestampWritesOnComputePass({{querySet, 1, wgpu::ComputePassTimestampLocation::End}}); + } + + // Set timestampWrites with same query set and query index on same compute pass + { + wgpu::QuerySet querySet = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnComputePass( + {{querySet, 0, wgpu::ComputePassTimestampLocation::Beginning}, + {querySet, 0, wgpu::ComputePassTimestampLocation::End}}); + } + + // Set timestampWrites with same query indexes and locations on different compute pass + { + wgpu::QuerySet querySet0 = CreateQuerySetForTimestamp(kQueryCount); + wgpu::QuerySet querySet1 = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnComputePass( + {{querySet0, 0, wgpu::ComputePassTimestampLocation::Beginning}, + {querySet0, 1, wgpu::ComputePassTimestampLocation::End}}, + {{querySet1, 0, wgpu::ComputePassTimestampLocation::Beginning}, + {querySet1, 1, wgpu::ComputePassTimestampLocation::End}}); + } +} + +// Test timestampWrites setting in render pass descriptor +TEST_P(TimestampQueryTests, TimestampWritesOnRenderPass) { + // TODO(dawn:1489): Fails on Intel Windows Vulkan due to a driver issue that + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until + // the issue is fixed. + DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel()); + + // TODO (dawn:1250): Still not implemented on Metal backend. + DAWN_TEST_UNSUPPORTED_IF(IsMetal()); + + constexpr uint32_t kQueryCount = 2; + + // Set timestampWrites with different query indexes and locations, not need test write same + // query index due to it's not allowed on render pass. + { + wgpu::QuerySet querySet = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnRenderPass( + {{querySet, 0, wgpu::RenderPassTimestampLocation::Beginning}, + {querySet, 1, wgpu::RenderPassTimestampLocation::End}}); + } + + // Set timestampWrites with different query set on same render pass + { + wgpu::QuerySet querySet0 = CreateQuerySetForTimestamp(1); + wgpu::QuerySet querySet1 = CreateQuerySetForTimestamp(1); + + TestTimestampWritesOnRenderPass( + {{querySet0, 0, wgpu::RenderPassTimestampLocation::Beginning}, + {querySet1, 0, wgpu::RenderPassTimestampLocation::End}}); + } + + // Set timestampWrites with only one value of RenderPassTimestampLocation + { + wgpu::QuerySet querySet = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnRenderPass( + {{querySet, 0, wgpu::RenderPassTimestampLocation::Beginning}}); + + TestTimestampWritesOnRenderPass({{querySet, 1, wgpu::RenderPassTimestampLocation::End}}); + } + + // Set timestampWrites with same query indexes and locations on different render pass + { + wgpu::QuerySet querySet0 = CreateQuerySetForTimestamp(kQueryCount); + wgpu::QuerySet querySet1 = CreateQuerySetForTimestamp(kQueryCount); + + TestTimestampWritesOnRenderPass( + {{querySet0, 0, wgpu::RenderPassTimestampLocation::Beginning}, + {querySet0, 1, wgpu::RenderPassTimestampLocation::End}}, + {{querySet1, 0, wgpu::RenderPassTimestampLocation::Beginning}, + {querySet1, 1, wgpu::RenderPassTimestampLocation::End}}); + } +} + // Test resolving timestamp query from another different encoder TEST_P(TimestampQueryTests, ResolveFromAnotherEncoder) { constexpr uint32_t kQueryCount = 2; @@ -717,8 +956,8 @@ TEST_P(TimestampQueryTests, ResolveFromAnotherEncoder) { // Test resolving timestamp query correctly if the queries are written sparsely TEST_P(TimestampQueryTests, ResolveSparseQueries) { - // TODO(hao.x.li@intel.com): Fails on Intel Windows Vulkan due to a driver issue that - // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it util + // TODO(dawn:1489): Fails on Intel Windows Vulkan due to a driver issue that + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until // the issue is fixed. DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel()); @@ -768,8 +1007,8 @@ TEST_P(TimestampQueryTests, ResolveWithoutWritten) { // Test resolving timestamp query to one slot in the buffer TEST_P(TimestampQueryTests, ResolveToBufferWithOffset) { - // TODO(hao.x.li@intel.com): Fails on Intel Windows Vulkan due to a driver issue that - // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it util + // TODO(dawn:1489): Fails on Intel Windows Vulkan due to a driver issue that + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until // the issue is fixed. DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel()); @@ -818,8 +1057,8 @@ TEST_P(TimestampQueryTests, ResolveToBufferWithOffset) { // Test resolving a query set twice into the same destination buffer with potentially overlapping // ranges TEST_P(TimestampQueryTests, ResolveTwiceToSameBuffer) { - // TODO(hao.x.li@intel.com): Fails on Intel Windows Vulkan due to a driver issue that - // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it util + // TODO(dawn:1489): Fails on Intel Windows Vulkan due to a driver issue that + // vkCmdFillBuffer and vkCmdCopyQueryPoolResults are not executed in order, skip it until // the issue is fixed. DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel());