diff --git a/pxr/imaging/hdSt/commandBuffer.cpp b/pxr/imaging/hdSt/commandBuffer.cpp index b2b7eac2b4..001a526db2 100644 --- a/pxr/imaging/hdSt/commandBuffer.cpp +++ b/pxr/imaging/hdSt/commandBuffer.cpp @@ -85,6 +85,8 @@ HdStCommandBuffer::PrepareDraw( for (auto const& batch : _drawBatches) { batch->PrepareDraw(gfxCmds, renderPassState, resourceRegistry); } + + resourceRegistry->SubmitComputeWork(); } void diff --git a/pxr/imaging/hdSt/pipelineDrawBatch.cpp b/pxr/imaging/hdSt/pipelineDrawBatch.cpp index 1d6645d4f9..53bd3be26a 100644 --- a/pxr/imaging/hdSt/pipelineDrawBatch.cpp +++ b/pxr/imaging/hdSt/pipelineDrawBatch.cpp @@ -48,6 +48,7 @@ #include "pxr/imaging/hgi/blitCmds.h" #include "pxr/imaging/hgi/blitCmdsOps.h" #include "pxr/imaging/hgi/graphicsPipeline.h" +#include "pxr/imaging/hgi/indirectCommandEncoder.h" #include "pxr/imaging/hgi/resourceBindings.h" #include "pxr/base/tf/diagnostic.h" @@ -931,7 +932,7 @@ HdSt_PipelineDrawBatch::_HasNothingToDraw() const void HdSt_PipelineDrawBatch::PrepareDraw( - HgiGraphicsCmds *, + HgiGraphicsCmds *gfxCmds, HdStRenderPassStateSharedPtr const & renderPassState, HdStResourceRegistrySharedPtr const & resourceRegistry) { @@ -953,6 +954,21 @@ HdSt_PipelineDrawBatch::PrepareDraw( _dispatchBuffer->CopyData(_drawCommandBuffer); _drawCommandBufferDirty = false; } + + Hgi *hgi = resourceRegistry->GetHgi(); + HgiCapabilities const *capabilities = hgi->GetCapabilities(); + + // For ICBs on Apple Silicon, we do not support rendering to non-MSAA + // surfaces, such as OIT as Volumetrics. Disable in these cases. + bool const drawICB = + capabilities->IsSet(HgiDeviceCapabilitiesBitsIndirectCommandBuffers) && + gfxCmds && + renderPassState->GetMultiSampleEnabled(); + + _indirectCommands.reset(); + if (drawICB) { + _PrepareIndirectCommandBuffer(renderPassState, resourceRegistry); + } if (_useGpuCulling) { // Ignore passed in gfxCmds for now since GPU frustum culling @@ -1235,56 +1251,68 @@ HdSt_PipelineDrawBatch::ExecuteDraw( if (!TF_VERIFY(_dispatchBuffer)) return; if (_HasNothingToDraw()) return; - - HgiCapabilities const *capabilities = - resourceRegistry->GetHgi()->GetCapabilities(); - - // Drawing can be either direct or indirect. For either case, - // the drawing batch and drawing program are prepared to resolve - // drawing coordinate state indirectly, i.e. from buffer data. - bool const drawIndirect = - capabilities->IsSet(HgiDeviceCapabilitiesBitsMultiDrawIndirect); - _DrawingProgram & program = _GetDrawingProgram(renderPassState, - resourceRegistry); - if (!TF_VERIFY(program.IsValid())) return; - - _BindingState state( - _drawItemInstances.front()->GetDrawItem(), - _dispatchBuffer, - program.GetBinder(), - program.GetGLSLProgram(), - program.GetComposedShaders(), - program.GetGeometricShader()); - - Hgi * hgi = resourceRegistry->GetHgi(); - - HgiGraphicsPipelineSharedPtr pso = - _GetDrawPipeline( - renderPassState, - resourceRegistry, - state); - HgiGraphicsPipelineHandle psoHandle = *pso.get(); - gfxCmds->BindPipeline(psoHandle); - - HgiResourceBindingsDesc bindingsDesc; - state.GetBindingsForDrawing(&bindingsDesc); - HgiResourceBindingsHandle resourceBindings = - hgi->CreateResourceBindings(bindingsDesc); - gfxCmds->BindResources(resourceBindings); + Hgi *hgi = resourceRegistry->GetHgi(); + HgiCapabilities const *capabilities = hgi->GetCapabilities(); + + bool const drawICB = + capabilities->IsSet(HgiDeviceCapabilitiesBitsIndirectCommandBuffers); - HgiVertexBufferBindingVector bindings; - _GetVertexBufferBindingsForDrawing(&bindings, state); - gfxCmds->BindVertexBuffers(bindings); + if (drawICB && _indirectCommands) { + HgiIndirectCommandEncoder *encoder = hgi->GetIndirectCommandEncoder(); + encoder->ExecuteDraw(gfxCmds, _indirectCommands.get()); - if (drawIndirect) { - _ExecuteDrawIndirect(gfxCmds, state.indexBar); - } else { - _ExecuteDrawImmediate(gfxCmds, state.indexBar); + hgi->DestroyResourceBindings(&(_indirectCommands->resourceBindings)); + _indirectCommands.reset(); } + else { + _DrawingProgram & program = _GetDrawingProgram(renderPassState, + resourceRegistry); + if (!TF_VERIFY(program.IsValid())) return; + + _BindingState state( + _drawItemInstances.front()->GetDrawItem(), + _dispatchBuffer, + program.GetBinder(), + program.GetGLSLProgram(), + program.GetComposedShaders(), + program.GetGeometricShader()); + + HgiGraphicsPipelineSharedPtr pso = + _GetDrawPipeline( + renderPassState, + resourceRegistry, + state); + + HgiGraphicsPipelineHandle psoHandle = *pso.get(); + gfxCmds->BindPipeline(psoHandle); + + HgiResourceBindingsDesc bindingsDesc; + state.GetBindingsForDrawing(&bindingsDesc); + + HgiResourceBindingsHandle resourceBindings = + hgi->CreateResourceBindings(bindingsDesc); + gfxCmds->BindResources(resourceBindings); + + HgiVertexBufferBindingVector bindings; + _GetVertexBufferBindingsForDrawing(&bindings, state); + gfxCmds->BindVertexBuffers(bindings); + + // Drawing can be either direct or indirect. For either case, + // the drawing batch and drawing program are prepared to resolve + // drawing coordinate state indirectly, i.e. from buffer data. + bool const drawIndirect = + capabilities->IsSet(HgiDeviceCapabilitiesBitsMultiDrawIndirect); + + if (drawIndirect) { + _ExecuteDrawIndirect(gfxCmds, state.indexBar); + } else { + _ExecuteDrawImmediate(gfxCmds, state.indexBar); + } - hgi->DestroyResourceBindings(&resourceBindings); + hgi->DestroyResourceBindings(&resourceBindings); + } HD_PERF_COUNTER_INCR(HdPerfTokens->drawCalls); HD_PERF_COUNTER_ADD(HdTokens->itemsDrawn, _numVisibleItems); @@ -1430,6 +1458,75 @@ _GetCullPipeline( return pipelineInstance.GetValue(); } +void +HdSt_PipelineDrawBatch::_PrepareIndirectCommandBuffer( + HdStRenderPassStateSharedPtr const & renderPassState, + HdStResourceRegistrySharedPtr const & resourceRegistry) +{ + Hgi *hgi = resourceRegistry->GetHgi(); + _DrawingProgram & program = _GetDrawingProgram(renderPassState, + resourceRegistry); + if (!TF_VERIFY(program.IsValid())) return; + + _BindingState state( + _drawItemInstances.front()->GetDrawItem(), + _dispatchBuffer, + program.GetBinder(), + program.GetGLSLProgram(), + program.GetComposedShaders(), + program.GetGeometricShader()); + + HgiGraphicsPipelineSharedPtr pso = + _GetDrawPipeline( + renderPassState, + resourceRegistry, + state); + + HgiGraphicsPipelineHandle psoHandle = *pso.get(); + + HgiResourceBindingsDesc bindingsDesc; + state.GetBindingsForDrawing(&bindingsDesc); + + HgiResourceBindingsHandle resourceBindings = + hgi->CreateResourceBindings(bindingsDesc); + + HgiVertexBufferBindingVector vertexBindings; + _GetVertexBufferBindingsForDrawing(&vertexBindings, state); + + HdStBufferResourceSharedPtr paramBuffer = _dispatchBuffer-> + GetBufferArrayRange()->GetResource(HdTokens->drawDispatch); + + HgiIndirectCommandEncoder *encoder = hgi->GetIndirectCommandEncoder(); + HgiComputeCmds *computeCmds = resourceRegistry->GetGlobalComputeCmds(); + + if (!_useDrawIndexed) { + _indirectCommands = encoder->EncodeDraw( + computeCmds, + psoHandle, + resourceBindings, + vertexBindings, + paramBuffer->GetHandle(), + paramBuffer->GetOffset(), + _dispatchBuffer->GetCount(), + paramBuffer->GetStride()); + } else { + HdStBufferResourceSharedPtr indexBuffer = + state.indexBar->GetResource(HdTokens->indices); + + _indirectCommands = encoder->EncodeDrawIndexed( + computeCmds, + psoHandle, + resourceBindings, + vertexBindings, + indexBuffer->GetHandle(), + paramBuffer->GetHandle(), + paramBuffer->GetOffset(), + _dispatchBuffer->GetCount(), + paramBuffer->GetStride(), + _patchBaseVertexByteOffset); + } +} + void HdSt_PipelineDrawBatch::_ExecuteFrustumCull( bool const updateBufferData, diff --git a/pxr/imaging/hdSt/pipelineDrawBatch.h b/pxr/imaging/hdSt/pipelineDrawBatch.h index 9bbfb7da12..f3ea96f826 100644 --- a/pxr/imaging/hdSt/pipelineDrawBatch.h +++ b/pxr/imaging/hdSt/pipelineDrawBatch.h @@ -35,6 +35,7 @@ PXR_NAMESPACE_OPEN_SCOPE class HgiCapabilities; +struct HgiIndirectCommands; using HdBindingRequestVector = std::vector; /// \class HdSt_PipelineDrawBatch @@ -128,6 +129,10 @@ class HdSt_PipelineDrawBatch : public HdSt_DrawBatch HdStResourceRegistrySharedPtr const & resourceRegistry); void _CompileBatch(HdStResourceRegistrySharedPtr const & resourceRegistry); + + void _PrepareIndirectCommandBuffer( + HdStRenderPassStateSharedPtr const & renderPassState, + HdStResourceRegistrySharedPtr const & resourceRegistry); bool _HasNothingToDraw() const; @@ -179,6 +184,8 @@ class HdSt_PipelineDrawBatch : public HdSt_DrawBatch size_t _instanceCountOffset; size_t _cullInstanceCountOffset; size_t _patchBaseVertexByteOffset; + + std::unique_ptr _indirectCommands; }; diff --git a/pxr/imaging/hgi/CMakeLists.txt b/pxr/imaging/hgi/CMakeLists.txt index 3e0257098d..53a45015ba 100644 --- a/pxr/imaging/hgi/CMakeLists.txt +++ b/pxr/imaging/hgi/CMakeLists.txt @@ -20,6 +20,7 @@ pxr_library(hgi graphicsCmdsDesc graphicsPipeline hgi + indirectCommandEncoder resourceBindings sampler shaderFunction diff --git a/pxr/imaging/hgi/enums.h b/pxr/imaging/hgi/enums.h index da9e396f3f..3ebb3c3bd9 100644 --- a/pxr/imaging/hgi/enums.h +++ b/pxr/imaging/hgi/enums.h @@ -72,6 +72,9 @@ using HgiBits = uint32_t; ///
  • HgiDeviceCapabilitiesBitsBasePrimitiveOffset: /// The device requires workaround for base primitive offset
    • /// +///
  • HgiDeviceCapabilitiesBitsIndirectCommandBuffers: +/// Indirect command buffers are supported
    • +/// /// enum HgiDeviceCapabilitiesBits : HgiBits { @@ -91,6 +94,7 @@ enum HgiDeviceCapabilitiesBits : HgiBits HgiDeviceCapabilitiesBitsCustomDepthRange = 1 << 13, HgiDeviceCapabilitiesBitsMetalTessellation = 1 << 14, HgiDeviceCapabilitiesBitsBasePrimitiveOffset = 1 << 15, + HgiDeviceCapabilitiesBitsIndirectCommandBuffers = 1 << 16, }; using HgiDeviceCapabilities = HgiBits; diff --git a/pxr/imaging/hgi/graphicsCmds.h b/pxr/imaging/hgi/graphicsCmds.h index 139fec7afe..4b66e97ebd 100644 --- a/pxr/imaging/hgi/graphicsCmds.h +++ b/pxr/imaging/hgi/graphicsCmds.h @@ -203,6 +203,7 @@ class HgiGraphicsCmds : public HgiCmds HGI_API virtual void MemoryBarrier(HgiMemoryBarrier barrier) = 0; + protected: HGI_API HgiGraphicsCmds(); diff --git a/pxr/imaging/hgi/hgi.h b/pxr/imaging/hgi/hgi.h index a7b476f5fb..9132f4d254 100644 --- a/pxr/imaging/hgi/hgi.h +++ b/pxr/imaging/hgi/hgi.h @@ -48,6 +48,7 @@ PXR_NAMESPACE_OPEN_SCOPE class HgiCapabilities; +class HgiIndirectCommandEncoder; using HgiUniquePtr = std::unique_ptr; @@ -295,6 +296,9 @@ class Hgi HGI_API virtual HgiCapabilities const* GetCapabilities() const = 0; + HGI_API + virtual HgiIndirectCommandEncoder* GetIndirectCommandEncoder() = 0; + /// Optionally called by client app at the start of a new rendering frame. /// We can't rely on StartFrame for anything important, because it is up to /// the external client to (optionally) call this and they may never do. diff --git a/pxr/imaging/hgi/indirectCommandEncoder.cpp b/pxr/imaging/hgi/indirectCommandEncoder.cpp new file mode 100644 index 0000000000..5378220152 --- /dev/null +++ b/pxr/imaging/hgi/indirectCommandEncoder.cpp @@ -0,0 +1,31 @@ +// +// Copyright 2022 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// + +#include "pxr/imaging/hgi/indirectCommandEncoder.h" + +PXR_NAMESPACE_OPEN_SCOPE + +HgiIndirectCommandEncoder::~HgiIndirectCommandEncoder() = default; + +PXR_NAMESPACE_CLOSE_SCOPE diff --git a/pxr/imaging/hgi/indirectCommandEncoder.h b/pxr/imaging/hgi/indirectCommandEncoder.h new file mode 100644 index 0000000000..f22e473dbc --- /dev/null +++ b/pxr/imaging/hgi/indirectCommandEncoder.h @@ -0,0 +1,112 @@ +// +// Copyright 2022 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// +#ifndef PXR_IMAGING_HGI_INDIRECT_COMMAND_ENCODER_H +#define PXR_IMAGING_HGI_INDIRECT_COMMAND_ENCODER_H + +#include "pxr/pxr.h" +#include "pxr/imaging/hgi/api.h" +#include "pxr/imaging/hgi/cmds.h" +#include "pxr/imaging/hgi/resourceBindings.h" +#include "pxr/imaging/hgi/graphicsPipeline.h" + +#include +#include + +PXR_NAMESPACE_OPEN_SCOPE + +class Hgi; +class HgiComputeCmds; +class HgiGraphicsCmds; + +struct HgiIndirectCommands +{ + HgiIndirectCommands(uint32_t drawCount, + HgiGraphicsPipelineHandle const &graphicsPipeline, + HgiResourceBindingsHandle const &resourceBindings) + : drawCount(drawCount) + , graphicsPipeline(graphicsPipeline) + , resourceBindings(resourceBindings) + { + } + + virtual ~HgiIndirectCommands() = default; + + uint32_t drawCount; + HgiGraphicsPipelineHandle graphicsPipeline; + HgiResourceBindingsHandle resourceBindings; +}; + +using HgiIndirectCommandsUniquePtr = std::unique_ptr; + +/// \class HgiIndirectCommandEncoder +/// +/// XXX: Explain what this is.... +/// +class HgiIndirectCommandEncoder : public HgiCmds +{ +public: + HGI_API + virtual ~HgiIndirectCommandEncoder(); + + HGI_API + virtual HgiIndirectCommandsUniquePtr EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride) = 0; + + HGI_API + virtual HgiIndirectCommandsUniquePtr EncodeDrawIndexed( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& indexBuffer, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) = 0; + + HGI_API + virtual void ExecuteDraw( + HgiGraphicsCmds * gfxCmds, + HgiIndirectCommands const* commands) = 0; + +protected: + HGI_API + HgiIndirectCommandEncoder() = default; + +private: + HgiIndirectCommandEncoder & operator=(const HgiIndirectCommandEncoder&) = delete; + HgiIndirectCommandEncoder(const HgiIndirectCommandEncoder&) = delete; +}; + +PXR_NAMESPACE_CLOSE_SCOPE + +#endif diff --git a/pxr/imaging/hgiGL/CMakeLists.txt b/pxr/imaging/hgiGL/CMakeLists.txt index 4776a6113c..3be1b0ee1c 100644 --- a/pxr/imaging/hgiGL/CMakeLists.txt +++ b/pxr/imaging/hgiGL/CMakeLists.txt @@ -35,6 +35,7 @@ pxr_library(hgiGL graphicsCmds graphicsPipeline hgi + indirectCommandEncoder ops resourceBindings sampler diff --git a/pxr/imaging/hgiGL/hgi.cpp b/pxr/imaging/hgiGL/hgi.cpp index d02de95864..c56e4a4edc 100644 --- a/pxr/imaging/hgiGL/hgi.cpp +++ b/pxr/imaging/hgiGL/hgi.cpp @@ -80,6 +80,7 @@ HgiGL::HgiGL() _device = new HgiGLDevice(); _capabilities.reset(new HgiGLCapabilities()); + _indirectCommandEncoder.reset(new HgiGLIndirectCommandEncoder()); } HgiGL::~HgiGL() @@ -260,6 +261,12 @@ HgiGL::GetCapabilities() const return _capabilities.get(); } +HgiGLIndirectCommandEncoder* +HgiGL::GetIndirectCommandEncoder() +{ + return _indirectCommandEncoder.get(); +} + void HgiGL::StartFrame() { diff --git a/pxr/imaging/hgiGL/hgi.h b/pxr/imaging/hgiGL/hgi.h index 9ad6e4a5d3..74669dbf5e 100644 --- a/pxr/imaging/hgiGL/hgi.h +++ b/pxr/imaging/hgiGL/hgi.h @@ -28,6 +28,7 @@ #include "pxr/imaging/hgiGL/api.h" #include "pxr/imaging/hgiGL/capabilities.h" #include "pxr/imaging/hgiGL/garbageCollector.h" +#include "pxr/imaging/hgiGL/indirectCommandEncoder.h" #include "pxr/imaging/hgi/hgi.h" #include "pxr/imaging/hgi/tokens.h" @@ -163,6 +164,9 @@ class HgiGL final : public Hgi HGIGL_API HgiGLCapabilities const* GetCapabilities() const override; + HGIGL_API + HgiGLIndirectCommandEncoder* GetIndirectCommandEncoder() override; + HGIGL_API void StartFrame() override; @@ -217,6 +221,7 @@ class HgiGL final : public Hgi HgiGLDevice* _device; std::unique_ptr _capabilities; + std::unique_ptr _indirectCommandEncoder; HgiGLGarbageCollector _garbageCollector; int _frameDepth; }; diff --git a/pxr/imaging/hgiGL/indirectCommandEncoder.cpp b/pxr/imaging/hgiGL/indirectCommandEncoder.cpp new file mode 100644 index 0000000000..98b5846f5b --- /dev/null +++ b/pxr/imaging/hgiGL/indirectCommandEncoder.cpp @@ -0,0 +1,70 @@ +// +// Copyright 2022 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// + +#include "pxr/imaging/hgiGL/indirectCommandEncoder.h" + +PXR_NAMESPACE_OPEN_SCOPE + +HgiGLIndirectCommandEncoder::HgiGLIndirectCommandEncoder() = default; + +HgiIndirectCommandsUniquePtr +HgiGLIndirectCommandEncoder::EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride) +{ + // No implementation + return nullptr; +} + +HgiIndirectCommandsUniquePtr +HgiGLIndirectCommandEncoder::EncodeDrawIndexed( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& indexBuffer, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) +{ + // No implementation + return nullptr; +} + +void +HgiGLIndirectCommandEncoder::ExecuteDraw( + HgiGraphicsCmds * gfxCmds, + HgiIndirectCommands const* commands) +{ + // No implementation +} +PXR_NAMESPACE_CLOSE_SCOPE diff --git a/pxr/imaging/hgiGL/indirectCommandEncoder.h b/pxr/imaging/hgiGL/indirectCommandEncoder.h new file mode 100644 index 0000000000..aa67a406ed --- /dev/null +++ b/pxr/imaging/hgiGL/indirectCommandEncoder.h @@ -0,0 +1,79 @@ +// +// Copyright 2022 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// +#ifndef PXR_IMAGING_HGI_GL_INDIRECT_COMMAND_ENCODER_H +#define PXR_IMAGING_HGI_GL_INDIRECT_COMMAND_ENCODER_H + +#include "pxr/pxr.h" +#include "pxr/imaging/hgiGL/api.h" +#include "pxr/imaging/hgi/indirectCommandEncoder.h" + +PXR_NAMESPACE_OPEN_SCOPE + +/// \class HgiGLIndirectGraphicsCmds +/// +/// Empty OpenGL implementation of Indirect Command Buffers. +/// +class HgiGLIndirectCommandEncoder final : public HgiIndirectCommandEncoder +{ +public: + HGIGL_API + HgiGLIndirectCommandEncoder(); + + HGIGL_API + HgiIndirectCommandsUniquePtr EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride) override; + + HGIGL_API + HgiIndirectCommandsUniquePtr EncodeDrawIndexed( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& indexBuffer, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) override; + + HGIGL_API + void ExecuteDraw( + HgiGraphicsCmds * gfxCmds, + HgiIndirectCommands const* commands) override; + +private: + HgiGLIndirectCommandEncoder & operator=(const HgiGLIndirectCommandEncoder&) = delete; + HgiGLIndirectCommandEncoder(const HgiGLIndirectCommandEncoder&) = delete; +}; + +PXR_NAMESPACE_CLOSE_SCOPE + +#endif diff --git a/pxr/imaging/hgiMetal/CMakeLists.txt b/pxr/imaging/hgiMetal/CMakeLists.txt index a817030f95..fef4a4fe0c 100644 --- a/pxr/imaging/hgiMetal/CMakeLists.txt +++ b/pxr/imaging/hgiMetal/CMakeLists.txt @@ -30,6 +30,7 @@ pxr_library(hgiMetal graphicsCmds.h graphicsPipeline.h hgi.h + indirectCommandEncoder.h resourceBindings.h sampler.h shaderFunction.h @@ -53,6 +54,7 @@ pxr_library(hgiMetal graphicsCmds.mm graphicsPipeline.mm hgi.mm + indirectCommandEncoder.mm resourceBindings.mm sampler.mm shaderFunction.mm diff --git a/pxr/imaging/hgiMetal/capabilities.mm b/pxr/imaging/hgiMetal/capabilities.mm index 529027317b..92fdfa1f10 100644 --- a/pxr/imaging/hgiMetal/capabilities.mm +++ b/pxr/imaging/hgiMetal/capabilities.mm @@ -39,6 +39,7 @@ bool unifiedMemory = false; bool barycentrics = false; bool hasAppleSilicon = false; + bool icbSupported = true; if (@available(macOS 100.100, ios 12.0, *)) { unifiedMemory = true; } else if (@available(macOS 10.15, ios 13.0, *)) { @@ -53,6 +54,15 @@ || [device areBarycentricCoordsSupported]; hasAppleSilicon = [device hasUnifiedMemory] && ![device isLowPower]; + + } + + // Indirect command buffers only on Apple Silicon GPUs with macOS 12.3 or later. + if (hasAppleSilicon) { + icbSupported = false; + if (@available(macOS 12.3, *)) { + icbSupported = true; + } } _SetFlag(HgiDeviceCapabilitiesBitsUnifiedMemory, unifiedMemory); @@ -68,6 +78,8 @@ _SetFlag(HgiDeviceCapabilitiesBitsMetalTessellation, true); _SetFlag(HgiDeviceCapabilitiesBitsMultiDrawIndirect, true); + + _SetFlag(HgiDeviceCapabilitiesBitsIndirectCommandBuffers, icbSupported); // This is done to decide whether to use a workaround for post tess // patch primitive ID lookup. The bug causes the firstPatch offset diff --git a/pxr/imaging/hgiMetal/computeCmds.h b/pxr/imaging/hgiMetal/computeCmds.h index 78bd591059..b75d9f40e5 100644 --- a/pxr/imaging/hgiMetal/computeCmds.h +++ b/pxr/imaging/hgiMetal/computeCmds.h @@ -71,6 +71,9 @@ class HgiMetalComputeCmds final : public HgiComputeCmds HGIMETAL_API void MemoryBarrier(HgiMemoryBarrier barrier) override; + HGIMETAL_API + id GetEncoder(); + protected: friend class HgiMetal; diff --git a/pxr/imaging/hgiMetal/computeCmds.mm b/pxr/imaging/hgiMetal/computeCmds.mm index 1b82a83aac..4c8725e825 100644 --- a/pxr/imaging/hgiMetal/computeCmds.mm +++ b/pxr/imaging/hgiMetal/computeCmds.mm @@ -55,7 +55,7 @@ HgiMetalComputeCmds::_CreateEncoder() { if (!_encoder) { - _commandBuffer = _hgi->GetPrimaryCommandBuffer(this); + _commandBuffer = _hgi->GetPrimaryCommandBuffer(this, false); if (_commandBuffer == nil) { _commandBuffer = _hgi->GetSecondaryCommandBuffer(); _secondaryCommandBuffer = true; @@ -209,4 +209,11 @@ return submittedWork; } +HGIMETAL_API +id HgiMetalComputeCmds::GetEncoder() +{ + _CreateEncoder(); + return _encoder; +} + PXR_NAMESPACE_CLOSE_SCOPE diff --git a/pxr/imaging/hgiMetal/graphicsCmds.h b/pxr/imaging/hgiMetal/graphicsCmds.h index e5742bb88c..b40d3c283a 100644 --- a/pxr/imaging/hgiMetal/graphicsCmds.h +++ b/pxr/imaging/hgiMetal/graphicsCmds.h @@ -27,7 +27,6 @@ #include "pxr/pxr.h" #include "pxr/base/gf/vec4i.h" #include "pxr/imaging/hgiMetal/api.h" -#include "pxr/imaging/hgiMetal/hgi.h" #include "pxr/imaging/hgiMetal/stepFunctions.h" #include "pxr/imaging/hgi/graphicsCmds.h" #include @@ -37,8 +36,8 @@ PXR_NAMESPACE_OPEN_SCOPE struct HgiGraphicsCmdsDesc; -class HgiMetalGraphicsPipeline; class HgiMetalResourceBindings; +class HgiMetalGraphicsPipeline; /// \class HgiMetalGraphicsCmds /// @@ -119,6 +118,9 @@ class HgiMetalGraphicsCmds final : public HgiGraphicsCmds HGIMETAL_API void EnableParallelEncoder(bool enable); + // Needs to be accessible from the Metal IndirectCommandEncoder + id GetEncoder(uint32_t encoderIndex = 0); + protected: friend class HgiMetal; @@ -136,7 +138,6 @@ class HgiMetalGraphicsCmds final : public HgiGraphicsCmds HgiMetalGraphicsCmds(const HgiMetalGraphicsCmds&) = delete; uint32_t _GetNumEncoders(); - id _GetEncoder(uint32_t encoderIndex = 0); void _SetNumberParallelEncoders(uint32_t numEncoders); void _SetCachedEncoderState(id encoder); mutable std::mutex _encoderLock; @@ -167,6 +168,7 @@ class HgiMetalGraphicsCmds final : public HgiGraphicsCmds HgiGraphicsCmdsDesc _descriptor; HgiPrimitiveType _primitiveType; uint32_t _primitiveIndexSize; + uint32_t _drawBufferBindingIndex; NSString* _debugLabel; bool _hasWork; bool _viewportSet; diff --git a/pxr/imaging/hgiMetal/graphicsCmds.mm b/pxr/imaging/hgiMetal/graphicsCmds.mm index b15bea4672..930e0f88f8 100644 --- a/pxr/imaging/hgiMetal/graphicsCmds.mm +++ b/pxr/imaging/hgiMetal/graphicsCmds.mm @@ -28,6 +28,7 @@ #include "pxr/imaging/hgiMetal/graphicsCmds.h" #include "pxr/imaging/hgiMetal/graphicsPipeline.h" #include "pxr/imaging/hgiMetal/hgi.h" +#include "pxr/imaging/hgiMetal/indirectCommandEncoder.h" #include "pxr/imaging/hgiMetal/resourceBindings.h" #include "pxr/imaging/hgiMetal/texture.h" @@ -82,6 +83,7 @@ , _descriptor(desc) , _primitiveType(HgiPrimitiveTypeTriangleList) , _primitiveIndexSize(0) + , _drawBufferBindingIndex(0) , _debugLabel(nil) , _hasWork(false) , _viewportSet(false) @@ -246,7 +248,7 @@ } if (hasClear) { - _GetEncoder(); + GetEncoder(); _CreateArgumentBuffer(); } } @@ -280,7 +282,7 @@ // Fix for Vega in macOS before 12.0. There is state leakage between // indirect draw of different prim types which results in a GPU crash. // Flush with a null draw through the direct path. - id encoder = _GetEncoder(); + id encoder = GetEncoder(); MTLPrimitiveType mtlType = HgiMetalConversions::GetPrimitiveType(_primitiveType); [encoder drawPrimitives:mtlType @@ -376,7 +378,7 @@ } id -HgiMetalGraphicsCmds::_GetEncoder(uint32_t encoderIndex) +HgiMetalGraphicsCmds::GetEncoder(uint32_t encoderIndex) { uint32_t numActiveEncoders = _GetNumEncoders(); @@ -536,7 +538,7 @@ _SyncArgumentBuffer(); MTLPrimitiveType type=HgiMetalConversions::GetPrimitiveType(_primitiveType); - id encoder = _GetEncoder(); + id encoder = GetEncoder(); _stepFunctions.SetVertexBufferOffsets(encoder, baseInstance); @@ -599,7 +601,7 @@ const uint32_t encoderCount = (i == numEncoders - 1) ? finalCount : normalCount; wd.Run([&, i, encoderOffset, encoderCount]() { - id encoder = _GetEncoder(i); + id encoder = GetEncoder(i); if (_primitiveType == HgiPrimitiveTypePatchList) { const NSUInteger controlPointCount = _primitiveIndexSize; @@ -652,7 +654,7 @@ MTLPrimitiveType mtlType = HgiMetalConversions::GetPrimitiveType(_primitiveType); - id encoder = _GetEncoder(); + id encoder = GetEncoder(); _stepFunctions.SetVertexBufferOffsets(encoder, baseInstance); @@ -723,7 +725,7 @@ const uint32_t encoderCount = (i == numEncoders - 1) ? finalCount : normalCount; wd.Run([&, i, encoderOffset, encoderCount]() { - id encoder = _GetEncoder(i); + id encoder = GetEncoder(i); if (_primitiveType == HgiPrimitiveTypePatchList) { const NSUInteger controlPointCount = _primitiveIndexSize; @@ -788,7 +790,7 @@ HGIMETAL_DEBUG_PUSH_GROUP(_parallelEncoder, label) } else if (!_encoders.empty()) { - HGIMETAL_DEBUG_PUSH_GROUP(_GetEncoder(), label) + HGIMETAL_DEBUG_PUSH_GROUP(GetEncoder(), label) } else { _debugLabel = [@(label) copy]; @@ -802,7 +804,7 @@ HGIMETAL_DEBUG_POP_GROUP(_parallelEncoder) } else if (!_encoders.empty()) { - HGIMETAL_DEBUG_POP_GROUP(_GetEncoder()); + HGIMETAL_DEBUG_POP_GROUP(GetEncoder()); } if (_debugLabel) { [_debugLabel release]; diff --git a/pxr/imaging/hgiMetal/graphicsPipeline.mm b/pxr/imaging/hgiMetal/graphicsPipeline.mm index 0cf392ec6f..cda8b9f4b2 100644 --- a/pxr/imaging/hgiMetal/graphicsPipeline.mm +++ b/pxr/imaging/hgiMetal/graphicsPipeline.mm @@ -23,6 +23,7 @@ // #include "pxr/imaging/hgiMetal/hgi.h" +#include "pxr/imaging/hgiMetal/capabilities.h" #include "pxr/imaging/hgiMetal/conversions.h" #include "pxr/imaging/hgiMetal/diagnostic.h" #include "pxr/imaging/hgiMetal/graphicsPipeline.h" @@ -124,6 +125,10 @@ // Create a new render pipeline state object HGIMETAL_DEBUG_LABEL(stateDesc, _descriptor.debugName.c_str()); stateDesc.rasterSampleCount = _descriptor.multiSampleState.sampleCount; + + bool const icbSupport = hgi->GetCapabilities()-> + IsSet(HgiDeviceCapabilitiesBitsIndirectCommandBuffers); + stateDesc.supportIndirectCommandBuffers = icbSupport; stateDesc.inputPrimitiveTopology = HgiMetalConversions::GetPrimitiveClass(_descriptor.primitiveType); diff --git a/pxr/imaging/hgiMetal/hgi.h b/pxr/imaging/hgiMetal/hgi.h index 1a4f31138e..e3a2a60b7f 100644 --- a/pxr/imaging/hgiMetal/hgi.h +++ b/pxr/imaging/hgiMetal/hgi.h @@ -27,6 +27,7 @@ #include "pxr/pxr.h" #include "pxr/imaging/hgiMetal/api.h" #include "pxr/imaging/hgiMetal/capabilities.h" +#include "pxr/imaging/hgiMetal/indirectCommandEncoder.h" #include "pxr/imaging/hgi/hgi.h" #include "pxr/imaging/hgi/tokens.h" @@ -143,6 +144,9 @@ class HgiMetal final : public Hgi HGIMETAL_API HgiMetalCapabilities const* GetCapabilities() const override; + HGIMETAL_API + HgiMetalIndirectCommandEncoder* GetIndirectCommandEncoder() override; + HGIMETAL_API void StartFrame() override; @@ -228,6 +232,7 @@ class HgiMetal final : public Hgi HgiCmds* _currentCmds; std::unique_ptr _capabilities; + std::unique_ptr _indirectCommandEncoder; int _frameDepth; bool _workToFlush; diff --git a/pxr/imaging/hgiMetal/hgi.mm b/pxr/imaging/hgiMetal/hgi.mm index 0f0abeb2cc..f5daeee4cf 100644 --- a/pxr/imaging/hgiMetal/hgi.mm +++ b/pxr/imaging/hgiMetal/hgi.mm @@ -83,6 +83,7 @@ [_commandBuffer retain]; _capabilities.reset(new HgiMetalCapabilities(_device)); + _indirectCommandEncoder.reset(new HgiMetalIndirectCommandEncoder(this)); MTLArgumentDescriptor *argumentDescBuffer = [[MTLArgumentDescriptor alloc] init]; @@ -331,6 +332,12 @@ return _capabilities.get(); } +HgiMetalIndirectCommandEncoder* +HgiMetal::GetIndirectCommandEncoder() +{ + return _indirectCommandEncoder.get(); +} + void HgiMetal::StartFrame() { @@ -400,8 +407,9 @@ } void -HgiMetal::CommitPrimaryCommandBuffer(CommitCommandBufferWaitType waitType, - bool forceNewBuffer) +HgiMetal::CommitPrimaryCommandBuffer( + CommitCommandBufferWaitType waitType, + bool forceNewBuffer) { if (!_workToFlush && !forceNewBuffer) { return; @@ -462,7 +470,8 @@ { std::lock_guard lock(_freeArgMutex); if (_freeArgBuffers.empty()) { - buffer = [_device newBufferWithLength:4096 options:options]; + buffer = [_device newBufferWithLength:HgiMetalArgumentOffsetSize + options:options]; } else { buffer = _freeArgBuffers.top(); diff --git a/pxr/imaging/hgiMetal/indirectCommandEncoder.h b/pxr/imaging/hgiMetal/indirectCommandEncoder.h new file mode 100644 index 0000000000..2b6ab6aa12 --- /dev/null +++ b/pxr/imaging/hgiMetal/indirectCommandEncoder.h @@ -0,0 +1,132 @@ +// +// Copyright 2022 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// +#ifndef PXR_IMAGING_HGI_METAL_INDIRECT_COMMAND_ENCODER_H +#define PXR_IMAGING_HGI_METAL_INDIRECT_COMMAND_ENCODER_H + +#include "pxr/pxr.h" +#include "pxr/imaging/hgi/indirectCommandEncoder.h" +#include "pxr/imaging/hgiMetal/api.h" +#include "pxr/imaging/hgiMetal/stepFunctions.h" + +#import + +#include +#include +#include +#include + +PXR_NAMESPACE_OPEN_SCOPE + +class Hgi; +struct HgiIndirectCommands; +class HgiMetal; +class HgiMetalGraphicsPipeline; +class HgiMetalResourceBindings; +class HgiMetalVertexBufferBindings; + +/// \class HgiMetalIndirectCommandEncoder +/// +/// Metal implementation of Indirect Command Buffers. +/// +class HgiMetalIndirectCommandEncoder final : public HgiIndirectCommandEncoder +{ +public: + HGIMETAL_API + HgiMetalIndirectCommandEncoder(Hgi *hgi); + + HGIMETAL_API + HgiIndirectCommandsUniquePtr EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride) override; + + HGIMETAL_API + HgiIndirectCommandsUniquePtr EncodeDrawIndexed( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& indexBuffer, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) override; + + HGIMETAL_API + void ExecuteDraw( + HgiGraphicsCmds * gfxCmds, + HgiIndirectCommands const* commands) override; +private: + HgiMetalIndirectCommandEncoder & operator=(const HgiMetalIndirectCommandEncoder&) = delete; + HgiMetalIndirectCommandEncoder(const HgiMetalIndirectCommandEncoder&) = delete; + + HGIMETAL_API + id _GetFunction(HgiGraphicsPipelineDesc const& pipelineDesc, + bool isIndexed); + + HGIMETAL_API + HgiIndirectCommandsUniquePtr _EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const &bindings, + HgiBufferHandle const &indexBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset); + + HGIMETAL_API + id _AllocateCommandBuffer(uint32_t drawCount); + + HGIMETAL_API + id _AllocateArgumentBuffer(); + + HgiMetal *_hgi; + id _device; + id _library; + id _argumentEncoder; + std::vector> _functions; + MTLResourceOptions _bufferStorageMode; + id _triangleTessFactors; + id _quadTessFactors; + + using FreeCommandBuffers = + std::multimap>; + using FreeArgumentBuffers = std::stack>; + + std::mutex _poolMutex; + FreeCommandBuffers _commandBufferPool; + FreeArgumentBuffers _argumentBufferPool; +}; + +PXR_NAMESPACE_CLOSE_SCOPE + +#endif diff --git a/pxr/imaging/hgiMetal/indirectCommandEncoder.mm b/pxr/imaging/hgiMetal/indirectCommandEncoder.mm new file mode 100644 index 0000000000..ed520c4d16 --- /dev/null +++ b/pxr/imaging/hgiMetal/indirectCommandEncoder.mm @@ -0,0 +1,796 @@ +// +// Copyright 2020 Pixar +// +// Licensed under the Apache License, Version 2.0 (the "Apache License") +// with the following modification; you may not use this file except in +// compliance with the Apache License and the following modification to it: +// Section 6. Trademarks. is deleted and replaced with: +// +// 6. Trademarks. This License does not grant permission to use the trade +// names, trademarks, service marks, or product names of the Licensor +// and its affiliates, except as required to comply with Section 4(c) of +// the License and to reproduce the content of the NOTICE file. +// +// You may obtain a copy of the Apache License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the Apache License with the above modification is +// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +// KIND, either express or implied. See the Apache License for the specific +// language governing permissions and limitations under the Apache License. +// +#include "pxr/base/arch/defines.h" + +#include "pxr/imaging/hgiMetal/indirectCommandEncoder.h" +#include "pxr/imaging/hgiMetal/buffer.h" +#include "pxr/imaging/hgiMetal/capabilities.h" +#include "pxr/imaging/hgiMetal/conversions.h" +#include "pxr/imaging/hgiMetal/computeCmds.h" +#include "pxr/imaging/hgiMetal/graphicsCmds.h" +#include "pxr/imaging/hgiMetal/graphicsPipeline.h" +#include "pxr/imaging/hgiMetal/hgi.h" +#include "pxr/imaging/hgiMetal/resourceBindings.h" +#include "pxr/imaging/hgiMetal/stepFunctions.h" + +#include + +PXR_NAMESPACE_OPEN_SCOPE + +static const uint32_t MaxVertexBufferBindings = 64; +static const uint32_t MaxStepFunctions = 16; +static const uint32_t StepFunctionSize = 3; +static const uint32_t StepFunctionsArraySize = MaxStepFunctions * StepFunctionSize; + +struct HgiMetalIndirectCommands : public HgiIndirectCommands +{ + HgiMetalIndirectCommands( + uint32_t drawCount, + HgiGraphicsPipelineHandle const &graphicsPipeline, + HgiResourceBindingsHandle const &resourceBindings, + id indirectCommandBuffer, + id argumentBuffer, + id mainArgumentBuffer) + : HgiIndirectCommands(drawCount, graphicsPipeline, resourceBindings) + , indirectCommandBuffer(indirectCommandBuffer) + , indirectArgumentBuffer(argumentBuffer) + , mainArgumentBuffer(mainArgumentBuffer) + { + } + + id indirectCommandBuffer; + id indirectArgumentBuffer; + id mainArgumentBuffer; +}; + +enum ArgIndex +{ + ArgIndex_ICB, + ArgIndex_IndexBuffer, + ArgIndex_MainArgumentBuffer, + ArgIndex_PatchFactorsBuffer, + ArgIndex_PrimitiveType, + ArgIndex_DrawStepFunction, + ArgIndex_NumPatchStepFuncs = ArgIndex_DrawStepFunction + StepFunctionSize, + ArgIndex_PatchStepFunctions, + ArgIndex_NumControlPoints = ArgIndex_PatchStepFunctions + StepFunctionsArraySize, + ArgIndex_PatchBaseVertexByteOffset, + ArgIndex_NumBuffers, + ArgIndex_Buffers, +}; + +HgiMetalIndirectCommandEncoder::HgiMetalIndirectCommandEncoder(Hgi* hgi) + : HgiIndirectCommandEncoder() + , _device(nil) + , _library(nil) + , _argumentEncoder(nil) + , _triangleTessFactors(nil) + , _quadTessFactors(nil) +{ + _hgi = static_cast(hgi); + _device = _hgi->GetPrimaryDevice(); + _bufferStorageMode = _hgi->GetCapabilities()->defaultStorageMode; + + uint16_t const factorZero = + reinterpret_cast(GfHalf(0.0f).bits()); + uint16_t const factorOne = + reinterpret_cast(GfHalf(1.0f).bits()); + + MTLTriangleTessellationFactorsHalf triangleFactors; + triangleFactors.insideTessellationFactor = factorZero; + triangleFactors.edgeTessellationFactor[0] = factorOne; + triangleFactors.edgeTessellationFactor[1] = factorOne; + triangleFactors.edgeTessellationFactor[2] = factorOne; + _triangleTessFactors = + [_device + newBufferWithBytes:&triangleFactors + length:sizeof(triangleFactors) + options:_bufferStorageMode]; + if (_bufferStorageMode != MTLStorageModeShared && + [_triangleTessFactors respondsToSelector:@selector(didModifyRange:)]) { + [_triangleTessFactors didModifyRange:{0, _triangleTessFactors.length}]; + } + + MTLQuadTessellationFactorsHalf quadFactors; + quadFactors.insideTessellationFactor[0] = factorZero; + quadFactors.insideTessellationFactor[1] = factorZero; + quadFactors.edgeTessellationFactor[0] = factorOne; + quadFactors.edgeTessellationFactor[1] = factorOne; + quadFactors.edgeTessellationFactor[2] = factorOne; + quadFactors.edgeTessellationFactor[3] = factorOne; + _quadTessFactors = + [_device + newBufferWithBytes:&quadFactors + length:sizeof(quadFactors) + options:_bufferStorageMode]; + if (_bufferStorageMode != MTLStorageModeShared && + [_quadTessFactors respondsToSelector:@selector(didModifyRange:)]) { + [_quadTessFactors didModifyRange:{0, _quadTessFactors.length}]; + } +} + +std::string +_ArgId(ArgIndex index) +{ + return "[[ id(" + std::to_string(index) + ") ]]"; +} + +std::string +_Buffer(uint32_t index) +{ + return "[[ buffer(" + std::to_string(index) + ") ]]"; +} + +std::string +_MainArgBuffer(uint32_t offset, uint32_t index) +{ + return "(const device uint8_t*)args->mainArgumentBuffer + " + + std::to_string(offset) + ", " + std::to_string(index); +} + +id +HgiMetalIndirectCommandEncoder::_GetFunction( + HgiGraphicsPipelineDesc const& pipelineDesc, + bool isIndexed) +{ + static const std::string _shaderSource = + "#include \n" + "using namespace metal;\n" + + "struct StepFunctionDesc\n" + "{\n" + " uint32_t bindingIndex;\n" + " uint32_t byteOffset;\n" + " uint32_t stride;\n" + "};\n" + + "struct ArgBuffer\n" + "{\n" + " command_buffer commandBuffer " + _ArgId(ArgIndex_ICB) + ";\n" + " const device uint32_t* indexBuffer " + _ArgId(ArgIndex_IndexBuffer) + ";\n" + " const device void* mainArgumentBuffer " + _ArgId(ArgIndex_MainArgumentBuffer) + ";\n" + " const device void* patchFactorsBuffer " + _ArgId(ArgIndex_PatchFactorsBuffer) + ";\n" + " primitive_type type " + _ArgId(ArgIndex_PrimitiveType) + ";\n" + " StepFunctionDesc drawStepFunc " + _ArgId(ArgIndex_DrawStepFunction) + ";\n" + " uint32_t numPatchStepFuncs " + _ArgId(ArgIndex_NumPatchStepFuncs) + ";\n" + " array patchStepFuncs " + _ArgId(ArgIndex_PatchStepFunctions) + ";\n" + " uint32_t numControlPoints " + _ArgId(ArgIndex_NumControlPoints) + ";\n" + " uint32_t patchBaseVertexByteOffset " + _ArgId(ArgIndex_PatchBaseVertexByteOffset) + ";\n" + " uint32_t numBuffers " + _ArgId(ArgIndex_NumBuffers) + ";\n" + " array buffers " + _ArgId(ArgIndex_Buffers) + ";\n" + "};\n" + + "void SetMainBuffers(\n" + " thread render_command& cmd,\n" + " const device ArgBuffer* args)\n" + "{\n" + " cmd.set_vertex_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetBufferVS, HgiMetalArgumentIndexBuffers) + ");\n" + " cmd.set_fragment_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetBufferFS, HgiMetalArgumentIndexBuffers) + ");\n" + " cmd.set_vertex_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetTextureVS, HgiMetalArgumentIndexTextures) + ");\n" + " cmd.set_fragment_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetTextureFS, HgiMetalArgumentIndexTextures) + ");\n" + " cmd.set_vertex_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetSamplerVS, HgiMetalArgumentIndexSamplers) + ");\n" + " cmd.set_fragment_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetSamplerFS, HgiMetalArgumentIndexSamplers) + ");\n" + " cmd.set_vertex_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetConstants, HgiMetalArgumentIndexConstants) + ");\n" + " cmd.set_fragment_buffer(" + _MainArgBuffer(HgiMetalArgumentOffsetConstants, HgiMetalArgumentIndexConstants) + ");\n" + "}\n" + + "const device uint8_t* GetDrawCmdPtr(\n" + " const device ArgBuffer* args,\n" + " uint drawItemIndex)\n" + "{\n" + " return\n" + " (const device uint8_t*)args->buffers[args->drawStepFunc.bindingIndex]" + " + args->drawStepFunc.byteOffset\n" + " + (drawItemIndex * args->drawStepFunc.stride);\n" + "}\n" + + "const device uint8_t* BufferOffset(\n" + " const device uint8_t* bufferPtr," + " const device StepFunctionDesc& stepFuncDesc," + " uint32_t base)\n" + "{\n" + " return bufferPtr + stepFuncDesc.byteOffset + (base * stepFuncDesc.stride);\n" + "}\n" + + "// From _DrawNonIndexedCommand\n" + "struct DrawPrimitivesCmd\n" + "{\n" + " uint vertexCount;\n" + " uint instanceCount;\n" + " uint vertexStart;\n" + " uint baseInstance;\n" + "};\n" + + "kernel void\n" + "MultiDrawPrimitives(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device uint8_t* drawCmdU8 = GetDrawCmdPtr(args, drawItemIndex);\n" + " const device DrawPrimitivesCmd* drawCmd = " + " (const device DrawPrimitivesCmd*)drawCmdU8;\n" + " for (uint32_t i = 0; i < args->numBuffers; ++i)\n" + " {\n" + " const device uint8_t* bufferPtr = (device uint8_t*)args->buffers[i];\n" + " if (i == args->drawStepFunc.bindingIndex) {\n" + " bufferPtr = BufferOffset(bufferPtr, args->drawStepFunc, drawCmd->baseInstance);\n" + " }\n" + " cmd.set_vertex_buffer(bufferPtr, i);\n" + " }\n" + " cmd.draw_primitives(\n" + " args->type,\n" + " drawCmd->vertexStart,\n" + " drawCmd->vertexCount,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseInstance);\n" + "}\n" + + "// From _DrawIndexedCommand\n" + "struct DrawIndexedCmd\n" + "{\n" + " uint32_t indexCount;\n" + " uint32_t instanceCount;\n" + " uint32_t indexStart;\n" + " int32_t baseVertex;\n" + " uint32_t baseInstance;\n" + "};\n" + + "kernel void\n" + "MultiDrawIndexedPrimitives(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device uint8_t* drawCmdU8 = GetDrawCmdPtr(args, drawItemIndex);\n" + " const device DrawIndexedCmd* drawCmd = " + " (const device DrawIndexedCmd*)drawCmdU8;\n" + " for (uint32_t i = 0; i < args->numBuffers; ++i)\n" + " {\n" + " const device uint8_t* bufferPtr = (device uint8_t*)args->buffers[i];\n" + " if (i == args->drawStepFunc.bindingIndex) {\n" + " bufferPtr = BufferOffset(bufferPtr, args->drawStepFunc, drawCmd->baseInstance);\n" + " }\n" + " cmd.set_vertex_buffer(bufferPtr, i);\n" + " }\n" + " cmd.draw_indexed_primitives(\n" + " args->type,\n" + " drawCmd->indexCount,\n" + " args->indexBuffer + drawCmd->indexStart,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseVertex,\n" + " drawCmd->baseInstance);\n" + "}\n" + + "// From MTLDrawPatchIndirectArguments\n" + "struct DrawPatchCmd\n" + "{\n" + " uint32_t patchCount;\n" + " uint32_t instanceCount;\n" + " uint32_t patchStart;\n" + " uint32_t baseInstance;\n" + "};\n" + + "const device DrawPatchCmd* SetPatchVertexBuffers(\n" + " thread render_command& cmd,\n" + " const device ArgBuffer* args,\n" + " uint drawItemIndex)\n" + "{\n" + " const device uint8_t* drawCmdU8 = GetDrawCmdPtr(args, drawItemIndex);\n" + " const device DrawPatchCmd* drawCmd = (const device DrawPatchCmd*)drawCmdU8;\n" + " uint32_t baseVertexIndex = \n" + " (args->patchBaseVertexByteOffset + drawItemIndex * args->drawStepFunc.stride) / sizeof(uint32_t);\n" + " const device uint32_t* drawCmdU32 = (const device uint32_t*)(args->buffers[args->drawStepFunc.bindingIndex]);\n" + " uint32_t baseVertex = drawCmdU32[baseVertexIndex];\n" + " for (uint32_t i = 0; i < args->numBuffers; ++i)\n" + " {\n" + " const device uint8_t* bufferPtr = (device uint8_t*)args->buffers[i];\n" + " if (i == args->drawStepFunc.bindingIndex) {\n" + " bufferPtr = BufferOffset(bufferPtr, args->drawStepFunc, drawCmd->baseInstance);\n" + " } else {\n" + " for (uint32_t j = 0; j < args->numPatchStepFuncs; ++j) {\n" + " if (args->patchStepFuncs[j].bindingIndex == i) {\n" + " bufferPtr = BufferOffset(bufferPtr, args->patchStepFuncs[j], baseVertex);\n" + " }\n" + " }\n" + " }\n" + " cmd.set_vertex_buffer(bufferPtr, i);\n" + " }\n" + " return drawCmd;\n" + "}\n" + + "kernel void\n" + "MultiDrawTriPatches(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device DrawPatchCmd* drawCmd =\n" + " SetPatchVertexBuffers(cmd, args, drawItemIndex);\n" + " cmd.draw_patches(\n" + " args->numControlPoints,\n" + " drawCmd->patchStart,\n" + " drawCmd->patchCount,\n" + " (const device uint*)nullptr,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseInstance,\n" + " (const device MTLTriangleTessellationFactorsHalf*)args->patchFactorsBuffer);\n" + "}\n" + + "kernel void\n" + "MultiDrawIndexedTriPatches(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device DrawPatchCmd* drawCmd =\n" + " SetPatchVertexBuffers(cmd, args, drawItemIndex);\n" + " cmd.draw_indexed_patches(\n" + " args->numControlPoints,\n" + " drawCmd->patchStart,\n" + " drawCmd->patchCount,\n" + " (const device uint*)nullptr,\n" + " args->indexBuffer,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseInstance,\n" + " (const device MTLTriangleTessellationFactorsHalf*)args->patchFactorsBuffer);\n" + "}\n" + + "kernel void\n" + "MultiDrawQuadPatches(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device DrawPatchCmd* drawCmd =\n" + " SetPatchVertexBuffers(cmd, args, drawItemIndex);\n" + " cmd.draw_patches(\n" + " args->numControlPoints,\n" + " drawCmd->patchStart,\n" + " drawCmd->patchCount,\n" + " (const device uint*)nullptr,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseInstance,\n" + " (const device MTLQuadTessellationFactorsHalf*)args->patchFactorsBuffer);\n" + "}\n" + + "kernel void\n" + "MultiDrawIndexedQuadPatches(uint drawItemIndex [[ thread_position_in_grid ]],\n" + "device ArgBuffer *args " + _Buffer(HgiMetalArgumentIndexICB) + ")\n" + "{\n" + " render_command cmd(args->commandBuffer, drawItemIndex);\n" + " SetMainBuffers(cmd, args);\n" + " const device DrawPatchCmd* drawCmd =\n" + " SetPatchVertexBuffers(cmd, args, drawItemIndex);\n" + " cmd.draw_indexed_patches(\n" + " args->numControlPoints,\n" + " drawCmd->patchStart,\n" + " drawCmd->patchCount,\n" + " (const device uint*)nullptr,\n" + " args->indexBuffer,\n" + " drawCmd->instanceCount,\n" + " drawCmd->baseInstance,\n" + " (const device MTLQuadTessellationFactorsHalf*)args->patchFactorsBuffer);\n" + "}\n" + ; + + if (!_library) { + NSError *error = NULL; + MTLCompileOptions *options = [[MTLCompileOptions alloc] init]; + options.fastMathEnabled = YES; + + NSString *shaderSource = + [NSString stringWithUTF8String:_shaderSource.c_str()]; + + _library = [_device newLibraryWithSource:shaderSource + options:options + error:&error]; + + _functions.resize(6); + + if (!_library) { + NSString *errStr = [error localizedDescription]; + TF_FATAL_CODING_ERROR( + "Failed to create multidraw pipeline state: %s", + [errStr UTF8String]); + } + } + + static NSString *const _functionNames[] = { + @"MultiDrawPrimitives", + @"MultiDrawIndexedPrimitives", + @"MultiDrawTriPatches", + @"MultiDrawIndexedTriPatches", + @"MultiDrawQuadPatches", + @"MultiDrawIndexedQuadPatches" + }; + + int32_t type = 0; + + if (pipelineDesc.primitiveType != HgiPrimitiveTypePatchList) { + type = (!isIndexed) ? 0 : 1; + } else { + if (pipelineDesc.tessellationState.patchType == + HgiTessellationState::PatchType::Triangle) { + type = (!isIndexed) ? 2 : 3; + } else { + type = (!isIndexed) ? 4 : 5; + } + } + + id function = _functions[type]; + + if (!function) { + function = [_library newFunctionWithName:_functionNames[type]]; + function.label = _functionNames[type]; + _functions[type] = function; + } + + if (!_argumentEncoder) { + _argumentEncoder = + [function newArgumentEncoderWithBufferIndex:HgiMetalArgumentIndexICB]; + } + + return function; +} + +static const uint32_t _RoundUpPow2(uint32_t x) +{ + if (x <= 1) return 1; + + x--; + x |= x >> 1; + x |= x >> 2; + x |= x >> 4; + x |= x >> 8; + x |= x >> 16; + x++; + + return x; +} + +id +HgiMetalIndirectCommandEncoder::_AllocateCommandBuffer(uint32_t drawCount) +{ + uint32_t bufferSize = _RoundUpPow2(drawCount); + id commandBuffer = nil; + + { + // Search for a buffer of the required size in the free pool. + std::lock_guard lock(_poolMutex); + FreeCommandBuffers::iterator it = _commandBufferPool.find(bufferSize); + + if (it != _commandBufferPool.end()) { + commandBuffer = it->second; + _commandBufferPool.erase(it); + } + } + + if (!commandBuffer) { + MTLIndirectCommandBufferDescriptor* descriptor = + [MTLIndirectCommandBufferDescriptor new]; + + descriptor.commandTypes = MTLIndirectCommandTypeDraw + | MTLIndirectCommandTypeDrawIndexed + | MTLIndirectCommandTypeDrawPatches + | MTLIndirectCommandTypeDrawIndexedPatches; + descriptor.inheritBuffers = NO; + descriptor.inheritPipelineState = YES; + descriptor.maxVertexBufferBindCount = 31; + descriptor.maxFragmentBufferBindCount = 31; + + commandBuffer = + [_device newIndirectCommandBufferWithDescriptor:descriptor + maxCommandCount:bufferSize + options:MTLResourceStorageModePrivate]; + } + + return commandBuffer; +} + +// XXX: the encodedLength doesn't change, so it doesn't need to be a param. +id +HgiMetalIndirectCommandEncoder::_AllocateArgumentBuffer() +{ + id buffer = nil; + + { + std::lock_guard lock(_poolMutex); + if (!_argumentBufferPool.empty()) { + buffer = _argumentBufferPool.top(); + _argumentBufferPool.pop(); + memset(buffer.contents, 0x00, buffer.length); + } + } + + if (!buffer) { + buffer = + [_device newBufferWithLength:_argumentEncoder.encodedLength + options:_bufferStorageMode]; + } + + return buffer; +} + +template +void _SetArg( + id argumentEncoder, + ArgIndex argumentIndex, + T value) +{ + *(T*)[argumentEncoder constantDataAtIndex:argumentIndex] = value; +} + +HgiIndirectCommandsUniquePtr +HgiMetalIndirectCommandEncoder::_EncodeDraw( + HgiComputeCmds *computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const &bindings, + HgiBufferHandle const &indexBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) +{ + + HgiGraphicsPipelineDesc const &pipelineDesc = pipeline->GetDescriptor(); + id function = _GetFunction(pipelineDesc, !!indexBuffer); + + // Create an argument buffer to hold all buffers and constants. + std::unique_ptr commands = + std::make_unique( + drawCount, + pipeline, + resourceBindings, + _AllocateCommandBuffer(drawCount), + _AllocateArgumentBuffer(), + _hgi->GetArgBuffer()); + + [_argumentEncoder setArgumentBuffer:commands->indirectArgumentBuffer + offset:0]; + + HgiMetalComputeCmds* metalComputeCmds = static_cast(computeCmds); + id encoder = metalComputeCmds->GetEncoder(); + + // Create the ICB and add it to the argument buffer. + [_argumentEncoder setIndirectCommandBuffer:commands->indirectCommandBuffer + atIndex:ArgIndex_ICB]; + [encoder useResource:commands->indirectCommandBuffer + usage:(MTLResourceUsageRead | MTLResourceUsageWrite)]; + + // Pass the main argument buffer through so the resources can be bound. + [_argumentEncoder setBuffer:commands->mainArgumentBuffer + offset:0 + atIndex:ArgIndex_MainArgumentBuffer]; + [encoder useResource:commands->mainArgumentBuffer + usage:(MTLResourceUsageRead | MTLResourceUsageWrite)]; + + // Add the constants to argument buffer. + MTLPrimitiveType mtlPrimitiveType = + HgiMetalConversions::GetPrimitiveType(pipelineDesc.primitiveType); + + HgiMetalStepFunctions stepFunctions(pipelineDesc, bindings); + + HgiMetalStepFunctionDesc drawCommandStep + {stepFunctions.GetDrawBufferIndex(), drawBufferByteOffset, stride}; + + TF_VERIFY(stepFunctions.GetPatchBaseDescs().size() <= MaxStepFunctions); + + uint32_t controlPointCount = + pipelineDesc.tessellationState.primitiveIndexSize; + + _SetArg(_argumentEncoder, ArgIndex_PrimitiveType, mtlPrimitiveType); + _SetArg(_argumentEncoder, ArgIndex_DrawStepFunction, drawCommandStep); + _SetArg(_argumentEncoder, ArgIndex_NumControlPoints, controlPointCount); + _SetArg(_argumentEncoder, ArgIndex_PatchBaseVertexByteOffset, + patchBaseVertexByteOffset); + _SetArg(_argumentEncoder, ArgIndex_NumBuffers, bindings.size()); + _SetArg(_argumentEncoder, ArgIndex_NumPatchStepFuncs, + (uint32_t)stepFunctions.GetPatchBaseDescs().size()); + HgiMetalStepFunctionDesc* argPatchStepDescs = (HgiMetalStepFunctionDesc*) + [_argumentEncoder constantDataAtIndex:ArgIndex_PatchStepFunctions]; + for (auto const& stepFuncDesc : stepFunctions.GetPatchBaseDescs()) { + *argPatchStepDescs = stepFuncDesc; + ++argPatchStepDescs; + } + + // If this is a patch primitive then add the constant tess factors. + if (pipelineDesc.primitiveType == HgiPrimitiveTypePatchList) { + id patchFactorsBuffer = _triangleTessFactors; + if (pipelineDesc.tessellationState.patchType == + HgiTessellationState::PatchType::Quad) { + patchFactorsBuffer = _quadTessFactors; + } + [_argumentEncoder setBuffer:patchFactorsBuffer + offset:0 + atIndex:ArgIndex_PatchFactorsBuffer]; + [encoder useResource:patchFactorsBuffer + usage:(MTLResourceUsageRead | MTLResourceUsageWrite)]; + } + + // Add the index buffer to the argument buffer. + if (indexBuffer) { + HgiMetalBuffer* mtlIndexBuffer = + static_cast(indexBuffer.Get()); + [_argumentEncoder setBuffer:mtlIndexBuffer->GetBufferId() + offset:0 + atIndex:ArgIndex_IndexBuffer]; + [encoder useResource:mtlIndexBuffer->GetBufferId() + usage:(MTLResourceUsageRead | MTLResourceUsageWrite)]; + } + + // Add the vertex buffers to the argument buffer so they can be bound. + TF_VERIFY(bindings.size() < MaxVertexBufferBindings); + uint32_t index = 0; + + for (auto const& binding : bindings) { + HgiMetalBuffer* mtlBuffer = + static_cast(binding.buffer.Get()); + [_argumentEncoder setBuffer:mtlBuffer->GetBufferId() + offset:binding.byteOffset + atIndex:ArgIndex_Buffers + index]; + [encoder useResource:mtlBuffer->GetBufferId() + usage:(MTLResourceUsageRead | MTLResourceUsageWrite)]; + index++; + } + + if (_bufferStorageMode != MTLStorageModeShared && + [commands->indirectArgumentBuffer + respondsToSelector:@selector(didModifyRange:)]) { + [commands->indirectArgumentBuffer + didModifyRange:{0, commands->indirectArgumentBuffer.length}]; + } + + // Create a compute pipeline state for the generation of the ICB. + NSError *error; + id icbPipelineState = + [_device newComputePipelineStateWithFunction:function + error:&error]; + + // Set pipeline state on the encoder and dispatch to populate the ICB + [encoder setComputePipelineState:icbPipelineState]; + [encoder setBuffer:commands->indirectArgumentBuffer + offset:0 + atIndex:HgiMetalArgumentIndexICB]; + + NSUInteger threadExecutionWidth = icbPipelineState.threadExecutionWidth; + + [encoder dispatchThreads:MTLSizeMake(drawCount, 1, 1) + threadsPerThreadgroup:MTLSizeMake(threadExecutionWidth, 1, 1)]; + + [icbPipelineState release]; + + return commands; +} + +void +HgiMetalIndirectCommandEncoder::ExecuteDraw( + HgiGraphicsCmds * gfxCmds, + HgiIndirectCommands const* commands) +{ + HgiMetalGraphicsCmds *metalGfxCmds = static_cast(gfxCmds); + id encoder = metalGfxCmds->GetEncoder(); + + id commandBuffer = + _hgi->GetPrimaryCommandBuffer(this, false); + HgiMetalIndirectCommands const* metalCommands = + static_cast(commands); + + // Bind the encoder pipeline and draw everything in the ICB + HgiMetalGraphicsPipeline* graphicsPipeline = + static_cast(metalCommands->graphicsPipeline.Get()); + graphicsPipeline->BindPipeline(encoder); + + // Bind the resources. + id mainArgumentBuffer = metalCommands->mainArgumentBuffer; + HgiMetalResourceBindings* resourceBindings = + static_cast(metalCommands->resourceBindings.Get()); + resourceBindings->BindResources(_hgi, + encoder, + mainArgumentBuffer); + + // Ensure the the main argument buffer is updated on managed hardware. + if (mainArgumentBuffer.storageMode != MTLStorageModeShared && + [mainArgumentBuffer respondsToSelector:@selector(didModifyRange:)]) { + + [mainArgumentBuffer didModifyRange:{0, mainArgumentBuffer.length}]; + } + + id indirectCommandBuffer = + metalCommands->indirectCommandBuffer; + id argumentBuffer = metalCommands->indirectArgumentBuffer; + [encoder setVertexBuffer:argumentBuffer + offset:0 + atIndex:HgiMetalArgumentIndexICB]; + + [encoder + executeCommandsInBuffer:indirectCommandBuffer + withRange:{0, metalCommands->drawCount}]; + + [commandBuffer addCompletedHandler:^(id cmdBuffer) + { + std::lock_guard lock(_poolMutex); + + _argumentBufferPool.push(argumentBuffer); + _commandBufferPool.insert({ indirectCommandBuffer.size, + indirectCommandBuffer }); + }]; +} + +HgiIndirectCommandsUniquePtr +HgiMetalIndirectCommandEncoder::EncodeDraw( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride) +{ + HgiIndirectCommandsUniquePtr commands = + _EncodeDraw( + computeCmds, + pipeline, + resourceBindings, + vertexBindings, + HgiBufferHandle(), + drawBufferByteOffset, + drawCount, + stride, + 0); + + return commands; +} + +HgiIndirectCommandsUniquePtr +HgiMetalIndirectCommandEncoder::EncodeDrawIndexed( + HgiComputeCmds * computeCmds, + HgiGraphicsPipelineHandle const& pipeline, + HgiResourceBindingsHandle const& resourceBindings, + HgiVertexBufferBindingVector const& vertexBindings, + HgiBufferHandle const& indexBuffer, + HgiBufferHandle const& drawParameterBuffer, + uint32_t drawBufferByteOffset, + uint32_t drawCount, + uint32_t stride, + uint32_t patchBaseVertexByteOffset) +{ + HgiIndirectCommandsUniquePtr commands = + _EncodeDraw( + computeCmds, + pipeline, + resourceBindings, + vertexBindings, + indexBuffer, + drawBufferByteOffset, + drawCount, + stride, + patchBaseVertexByteOffset); + + return commands; +} + +PXR_NAMESPACE_CLOSE_SCOPE diff --git a/pxr/imaging/hgiMetal/resourceBindings.h b/pxr/imaging/hgiMetal/resourceBindings.h index a800cda3ce..8111121354 100644 --- a/pxr/imaging/hgiMetal/resourceBindings.h +++ b/pxr/imaging/hgiMetal/resourceBindings.h @@ -34,6 +34,7 @@ PXR_NAMESPACE_OPEN_SCOPE /// Chosen to be at the top of the range of indexs to not interfere /// with the vertex attributes. enum HgiMetalArgumentIndex { + HgiMetalArgumentIndexICB = 26, HgiMetalArgumentIndexConstants = 27, HgiMetalArgumentIndexSamplers = 28, HgiMetalArgumentIndexTextures = 29, diff --git a/pxr/imaging/hgiMetal/resourceBindings.mm b/pxr/imaging/hgiMetal/resourceBindings.mm index 20a78eb7d6..617967d4cc 100644 --- a/pxr/imaging/hgiMetal/resourceBindings.mm +++ b/pxr/imaging/hgiMetal/resourceBindings.mm @@ -45,7 +45,6 @@ id renderEncoder, id argBuffer) { - id device = hgi->GetPrimaryDevice(); id argEncoderBuffer = hgi->GetBufferArgumentEncoder(); id argEncoderSampler = hgi->GetSamplerArgumentEncoder(); id argEncoderTexture = hgi->GetTextureArgumentEncoder(); @@ -202,7 +201,6 @@ id computeEncoder, id argBuffer) { - id device = hgi->GetPrimaryDevice(); id argEncoderBuffer = hgi->GetBufferArgumentEncoder(); id argEncoderSampler = hgi->GetSamplerArgumentEncoder(); id argEncoderTexture = hgi->GetTextureArgumentEncoder();