diff options
45 files changed, 2607 insertions, 811 deletions
diff --git a/source/compiler-core/slang-metal-compiler.cpp b/source/compiler-core/slang-metal-compiler.cpp index d3bf92bf3..e235db23d 100644 --- a/source/compiler-core/slang-metal-compiler.cpp +++ b/source/compiler-core/slang-metal-compiler.cpp @@ -62,11 +62,33 @@ namespace Slang static SlangResult locateMetalCompiler(const String& path, DownstreamCompilerSet* set) { ComPtr<IDownstreamCompiler> innerCppCompiler; - ExecutableLocation exeLocation(path, "metal"); - SLANG_RETURN_ON_FAIL(GCCDownstreamCompilerUtil::createCompiler(exeLocation, innerCppCompiler)); + + ExecutableLocation metalcLocation = ExecutableLocation(path, "metal");; + + String metalSDKPath = path; + +#if defined (SLANG_APPLE_FAMILY) + // Use xcrun command to find the metal compiler. + CommandLine xcrunCmdLine; + ExecutableLocation xcrunLocation("xcrun"); + xcrunCmdLine.setExecutableLocation(xcrunLocation); + xcrunCmdLine.addArg("--sdk"); + xcrunCmdLine.addArg("macosx"); + xcrunCmdLine.addArg("--find"); + xcrunCmdLine.addArg("metal"); + ExecuteResult exeRes; + if (SLANG_SUCCEEDED(ProcessUtil::execute(xcrunCmdLine, exeRes))) + { + String metalPath = exeRes.standardOutput.trim(); + metalcLocation = ExecutableLocation(ExecutableLocation::Type::Path, metalPath); + metalSDKPath = Path::getParentDirectory(metalcLocation.m_pathOrName); + } +#endif + + SLANG_RETURN_ON_FAIL(GCCDownstreamCompilerUtil::createCompiler(metalcLocation, innerCppCompiler)); ComPtr<IDownstreamCompiler> compiler = ComPtr<IDownstreamCompiler>( - new MetalDownstreamCompiler(innerCppCompiler, path)); + new MetalDownstreamCompiler(innerCppCompiler, metalSDKPath)); set->addCompiler(compiler); return SLANG_OK; } diff --git a/source/slang/slang-reflection-api.cpp b/source/slang/slang-reflection-api.cpp index ab73ce7f4..e7f4b9bf3 100644 --- a/source/slang/slang-reflection-api.cpp +++ b/source/slang/slang-reflection-api.cpp @@ -1828,6 +1828,7 @@ namespace Slang case LayoutResourceKind::ShaderResource: case LayoutResourceKind::DescriptorTableSlot: case LayoutResourceKind::Uniform: + case LayoutResourceKind::ConstantBuffer: // for metal resInfo = info; break; } diff --git a/tools/gfx/metal/metal-base.h b/tools/gfx/metal/metal-base.h index 94577f145..b1650a655 100644 --- a/tools/gfx/metal/metal-base.h +++ b/tools/gfx/metal/metal-base.h @@ -21,6 +21,7 @@ namespace metal class TextureResourceImpl; class SamplerStateImpl; class ResourceViewImpl; + class BufferResourceViewImpl; class TextureResourceViewImpl; class TexelBufferResourceViewImpl; class PlainBufferResourceViewImpl; @@ -32,12 +33,12 @@ namespace metal class RayTracingPipelineStateImpl; class ShaderObjectLayoutImpl; class EntryPointLayout; - class RootShaderObjectLayout; + class RootShaderObjectLayoutImpl; class ShaderProgramImpl; class PipelineCommandEncoder; class ShaderObjectImpl; class MutableShaderObjectImpl; - //class RootShaderObjectImpl; + class RootShaderObjectImpl; class ShaderTableImpl; class ResourceCommandEncoder; class RenderCommandEncoder; diff --git a/tools/gfx/metal/metal-buffer.cpp b/tools/gfx/metal/metal-buffer.cpp index de866bcf7..4f2964c5b 100644 --- a/tools/gfx/metal/metal-buffer.cpp +++ b/tools/gfx/metal/metal-buffer.cpp @@ -1,5 +1,6 @@ // metal-buffer.cpp #include "metal-buffer.h" +#include "metal-util.h" namespace gfx { @@ -9,48 +10,49 @@ using namespace Slang; namespace metal { -BufferResourceImpl::BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* renderer) +BufferResourceImpl::BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* device) : Parent(desc) - , m_renderer(renderer) + , m_device(device) { - assert(renderer); } BufferResourceImpl::~BufferResourceImpl() { - if (sharedHandle.handleValue != 0) - { - } } DeviceAddress BufferResourceImpl::getDeviceAddress() { - return (DeviceAddress)0; + return m_buffer->gpuAddress(); } Result BufferResourceImpl::getNativeResourceHandle(InteropHandle* outHandle) { - return SLANG_E_NOT_IMPLEMENTED; + outHandle->api = InteropHandleAPI::Metal; + outHandle->handleValue = reinterpret_cast<intptr_t>(m_buffer.get()); + return SLANG_OK; } Result BufferResourceImpl::getSharedHandle(InteropHandle* outHandle) { - return SLANG_E_NOT_IMPLEMENTED; + return SLANG_E_NOT_AVAILABLE; } Result BufferResourceImpl::map(MemoryRange* rangeToRead, void** outPointer) { - return SLANG_E_NOT_IMPLEMENTED; + *outPointer = m_buffer->contents(); + return SLANG_OK; } Result BufferResourceImpl::unmap(MemoryRange* writtenRange) { - return SLANG_E_NOT_IMPLEMENTED; + return SLANG_OK; } Result BufferResourceImpl::setDebugName(const char* name) { - return SLANG_E_NOT_IMPLEMENTED; + Parent::setDebugName(name); + m_buffer->addDebugMarker(MetalUtil::createString(name).get(), NS::Range(0, m_desc.sizeInBytes)); + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-buffer.h b/tools/gfx/metal/metal-buffer.h index 96a3538f2..08fa6481d 100644 --- a/tools/gfx/metal/metal-buffer.h +++ b/tools/gfx/metal/metal-buffer.h @@ -17,22 +17,19 @@ class BufferResourceImpl : public BufferResource public: typedef BufferResource Parent; - BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* renderer); + RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::Buffer> m_buffer; + BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* device); ~BufferResourceImpl(); - DeviceImpl* m_renderer; - MTL::Buffer* m_buffer = nullptr; - virtual SLANG_NO_THROW DeviceAddress SLANG_MCALL getDeviceAddress() override; - virtual SLANG_NO_THROW Result SLANG_MCALL - getNativeResourceHandle(InteropHandle* outHandle) override; + virtual SLANG_NO_THROW Result SLANG_MCALL getNativeResourceHandle(InteropHandle* outHandle) override; virtual SLANG_NO_THROW Result SLANG_MCALL getSharedHandle(InteropHandle* outHandle) override; - virtual SLANG_NO_THROW Result SLANG_MCALL - map(MemoryRange* rangeToRead, void** outPointer) override; + virtual SLANG_NO_THROW Result SLANG_MCALL map(MemoryRange* rangeToRead, void** outPointer) override; virtual SLANG_NO_THROW Result SLANG_MCALL unmap(MemoryRange* writtenRange) override; diff --git a/tools/gfx/metal/metal-command-buffer.cpp b/tools/gfx/metal/metal-command-buffer.cpp index 08caa03dd..8aac6ea4b 100644 --- a/tools/gfx/metal/metal-command-buffer.cpp +++ b/tools/gfx/metal/metal-command-buffer.cpp @@ -21,12 +21,10 @@ ICommandBuffer* CommandBufferImpl::getInterface(const Guid& guid) return nullptr; } -void CommandBufferImpl::comFree() { } - -Result CommandBufferImpl::init(DeviceImpl* renderer, TransientResourceHeapImpl* transientHeap) +Result CommandBufferImpl::init(DeviceImpl* device, TransientResourceHeapImpl* transientHeap) { - m_renderer = renderer; - m_commandBuffer = m_renderer->m_commandQueue->commandBuffer(); + m_device = device; + m_commandBuffer = NS::RetainPtr(m_device->m_commandQueue->commandBuffer()); return SLANG_OK; } @@ -82,5 +80,55 @@ Result CommandBufferImpl::getNativeHandle(InteropHandle* outHandle) return SLANG_E_NOT_IMPLEMENTED; } +MTL::RenderCommandEncoder* CommandBufferImpl::getMetalRenderCommandEncoder() +{ + if (!m_metalRenderCommandEncoder) + { + endMetalCommandEncoder(); + // m_metalRenderCommandEncoder = NS::RetainPtr(m_commandBuffer->renderCommandEncoder()); + } + return m_metalRenderCommandEncoder.get(); +} + +MTL::ComputeCommandEncoder* CommandBufferImpl::getMetalComputeCommandEncoder() +{ + if (!m_metalComputeCommandEncoder) + { + endMetalCommandEncoder(); + m_metalComputeCommandEncoder = NS::RetainPtr(m_commandBuffer->computeCommandEncoder()); + } + return m_metalComputeCommandEncoder.get(); +} + +MTL::BlitCommandEncoder* CommandBufferImpl::getMetalBlitCommandEncoder() +{ + if (!m_metalBlitCommandEncoder) + { + endMetalCommandEncoder(); + m_metalBlitCommandEncoder = NS::RetainPtr(m_commandBuffer->blitCommandEncoder()); + } + return m_metalBlitCommandEncoder.get(); +} + +void CommandBufferImpl::endMetalCommandEncoder() +{ + if (m_metalRenderCommandEncoder) + { + m_metalRenderCommandEncoder->endEncoding(); + m_metalRenderCommandEncoder.reset(); + } + if (m_metalComputeCommandEncoder) + { + m_metalComputeCommandEncoder->endEncoding(); + m_metalComputeCommandEncoder.reset(); + } + if (m_metalBlitCommandEncoder) + { + m_metalBlitCommandEncoder->endEncoding(); + m_metalBlitCommandEncoder.reset(); + } +} + + } // namespace metal } // namespace gfx diff --git a/tools/gfx/metal/metal-command-buffer.h b/tools/gfx/metal/metal-command-buffer.h index aef616036..3041ba310 100644 --- a/tools/gfx/metal/metal-command-buffer.h +++ b/tools/gfx/metal/metal-command-buffer.h @@ -21,28 +21,35 @@ class CommandBufferImpl public: SLANG_COM_OBJECT_IUNKNOWN_ALL ICommandBuffer* getInterface(const Guid& guid); - virtual void comFree() override; public: - MTL::CommandBuffer* m_commandBuffer = nullptr; - DeviceImpl* m_renderer; - //bool m_isPreCommandBufferEmpty = true; + RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::CommandBuffer> m_commandBuffer; RootShaderObjectImpl m_rootObject; + // RefPtr<MutableRootShaderObjectImpl> m_mutableRootShaderObject; ResourceCommandEncoder* m_resourceCommandEncoder = nullptr; ComputeCommandEncoder* m_computeCommandEncoder = nullptr; RenderCommandEncoder* m_renderCommandEncoder = nullptr; RayTracingCommandEncoder* m_rayTracingCommandEncoder = nullptr; + NS::SharedPtr<MTL::RenderCommandEncoder> m_metalRenderCommandEncoder; + NS::SharedPtr<MTL::ComputeCommandEncoder> m_metalComputeCommandEncoder; + NS::SharedPtr<MTL::BlitCommandEncoder> m_metalBlitCommandEncoder; + // Command buffers are deallocated by its command pool, // so no need to free individually. ~CommandBufferImpl() = default; - using TransientResourceHeapImpl = gfx::SimpleTransientResourceHeap<DeviceImpl, CommandBufferImpl>; - Result init(DeviceImpl* renderer, TransientResourceHeapImpl* transientHeap); + Result init(DeviceImpl* device, TransientResourceHeapImpl* transientHeap); void beginCommandBuffer(); + MTL::RenderCommandEncoder* getMetalRenderCommandEncoder(); + MTL::ComputeCommandEncoder* getMetalComputeCommandEncoder(); + MTL::BlitCommandEncoder* getMetalBlitCommandEncoder(); + void endMetalCommandEncoder(); + public: virtual SLANG_NO_THROW void SLANG_MCALL encodeRenderCommands( IRenderPassLayout* renderPass, diff --git a/tools/gfx/metal/metal-command-encoder.cpp b/tools/gfx/metal/metal-command-encoder.cpp index 1e74733d1..f1e834400 100644 --- a/tools/gfx/metal/metal-command-encoder.cpp +++ b/tools/gfx/metal/metal-command-encoder.cpp @@ -10,6 +10,7 @@ #include "metal-shader-program.h" #include "metal-shader-table.h" #include "metal-texture.h" +#include "metal-util.h" #include "metal-helper-functions.h" @@ -24,37 +25,46 @@ namespace metal void PipelineCommandEncoder::init(CommandBufferImpl* commandBuffer) { m_commandBuffer = commandBuffer; - m_device = commandBuffer->m_renderer; - m_metalCommandBuffer = m_commandBuffer->m_commandBuffer; + m_metalCommandBuffer = m_commandBuffer->m_commandBuffer.get(); } -void ResourceCommandEncoder::copyBuffer( - IBufferResource* dst, Offset dstOffset, IBufferResource* src, Offset srcOffset, Size size) +void PipelineCommandEncoder::endEncodingImpl() { + m_commandBuffer->endMetalCommandEncoder(); } -void ResourceCommandEncoder::uploadBufferData( - IBufferResource* buffer, Offset offset, Size size, void* data) -{ -} - -void ResourceCommandEncoder::textureBarrier( - GfxCount count, ITextureResource* const* textures, ResourceState src, ResourceState dst) +Result PipelineCommandEncoder::setPipelineStateImpl(IPipelineState* state, IShaderObject** outRootObject) { + m_currentPipeline = static_cast<PipelineStateImpl*>(state); + // m_commandBuffer->m_mutableRootShaderObject = nullptr; + SLANG_RETURN_ON_FAIL(m_commandBuffer->m_rootObject.init( + m_commandBuffer->m_device, + m_currentPipeline->getProgram<ShaderProgramImpl>()->m_rootObjectLayout)); + *outRootObject = &m_commandBuffer->m_rootObject; + return SLANG_OK; } -// TODO: Change size_t to Count? -void ResourceCommandEncoder::bufferBarrier( - GfxCount count, IBufferResource* const* buffers, ResourceState src, ResourceState dst) +void ResourceCommandEncoder::endEncoding() { + PipelineCommandEncoder::endEncodingImpl(); } -void ResourceCommandEncoder::endEncoding() +void ResourceCommandEncoder::writeTimestamp(IQueryPool* queryPool, GfxIndex index) { + auto encoder = m_commandBuffer->getMetalBlitCommandEncoder(); + encoder->sampleCountersInBuffer(static_cast<QueryPoolImpl*>(queryPool)->m_counterSampleBuffer.get(), index, true); } -void ResourceCommandEncoder::writeTimestamp(IQueryPool* queryPool, GfxIndex index) +void ResourceCommandEncoder::copyBuffer( + IBufferResource* dst, Offset dstOffset, IBufferResource* src, Offset srcOffset, Size size) { + auto encoder = m_commandBuffer->getMetalBlitCommandEncoder(); + encoder->copyFromBuffer( + static_cast<BufferResourceImpl*>(src)->m_buffer.get(), + srcOffset, + static_cast<BufferResourceImpl*>(dst)->m_buffer.get(), + dstOffset, + size); } void ResourceCommandEncoder::copyTexture( @@ -68,6 +78,62 @@ void ResourceCommandEncoder::copyTexture( ITextureResource::Offset3D srcOffset, ITextureResource::Extents extent) { + auto encoder = m_commandBuffer->getMetalBlitCommandEncoder(); + + if (dstSubresource.layerCount == 0 && dstSubresource.mipLevelCount == 0 && srcSubresource.layerCount == 0 && srcSubresource.mipLevelCount == 0) + { + encoder->copyFromTexture( + static_cast<TextureResourceImpl*>(src)->m_texture.get(), + static_cast<TextureResourceImpl*>(dst)->m_texture.get()); + } + else + { + for (GfxIndex layer = 0; layer < dstSubresource.layerCount; layer++) + { + encoder->copyFromTexture( + static_cast<TextureResourceImpl*>(src)->m_texture.get(), + srcSubresource.baseArrayLayer + layer, + srcSubresource.mipLevel, + MTL::Origin(srcOffset.x, srcOffset.y, srcOffset.z), + MTL::Size(extent.width, extent.height, extent.depth), + static_cast<TextureResourceImpl*>(dst)->m_texture.get(), + dstSubresource.baseArrayLayer + layer, + dstSubresource.mipLevel, + MTL::Origin(dstOffset.x, dstOffset.y, dstOffset.z)); + } + } +} + +void ResourceCommandEncoder::copyTextureToBuffer( + IBufferResource* dst, + Offset dstOffset, + Size dstSize, + Size dstRowStride, + ITextureResource* src, + ResourceState srcState, + SubresourceRange srcSubresource, + ITextureResource::Offset3D srcOffset, + ITextureResource::Extents extent) +{ + assert(srcSubresource.mipLevelCount <= 1); + + auto encoder = m_commandBuffer->getMetalBlitCommandEncoder(); + encoder->copyFromTexture( + static_cast<TextureResourceImpl*>(src)->m_texture.get(), + srcSubresource.baseArrayLayer, + srcSubresource.mipLevel, + MTL::Origin(srcOffset.x, srcOffset.y, srcOffset.z), + MTL::Size(extent.width, extent.height, extent.depth), + static_cast<BufferResourceImpl*>(dst)->m_buffer.get(), + dstOffset, + dstRowStride, + dstSize); +} + +void ResourceCommandEncoder::uploadBufferData( + IBufferResource* buffer, Offset offset, Size size, void* data) +{ + SLANG_UNIMPLEMENTED_X("uploadBufferData"); } void ResourceCommandEncoder::uploadTextureData( @@ -78,12 +144,34 @@ void ResourceCommandEncoder::uploadTextureData( ITextureResource::SubresourceData* subResourceData, GfxCount subResourceDataCount) { + SLANG_UNIMPLEMENTED_X("uploadTextureData"); } +void ResourceCommandEncoder::bufferBarrier( + GfxCount count, IBufferResource* const* buffers, ResourceState src, ResourceState dst) +{ + // We use automatic hazard tracking for now, no need for barriers. +} + +void ResourceCommandEncoder::textureBarrier( + GfxCount count, ITextureResource* const* textures, ResourceState src, ResourceState dst) +{ + // We use automatic hazard tracking for now, no need for barriers. +} + +void ResourceCommandEncoder::textureSubresourceBarrier( + ITextureResource* texture, + SubresourceRange subresourceRange, + ResourceState src, + ResourceState dst) +{ + // We use automatic hazard tracking for now, no need for barriers. +} void ResourceCommandEncoder::clearResourceView( IResourceView* view, ClearValue* clearValue, ClearResourceViewFlags::Enum flags) { + SLANG_UNIMPLEMENTED_X("clearResourceView"); } void ResourceCommandEncoder::resolveResource( @@ -94,41 +182,29 @@ void ResourceCommandEncoder::resolveResource( ResourceState destState, SubresourceRange destRange) { + SLANG_UNIMPLEMENTED_X("resolveResource"); } void ResourceCommandEncoder::resolveQuery( IQueryPool* queryPool, GfxIndex index, GfxCount count, IBufferResource* buffer, Offset offset) { -} - -void ResourceCommandEncoder::copyTextureToBuffer( - IBufferResource* dst, - Offset dstOffset, - Size dstSize, - Size dstRowStride, - ITextureResource* src, - ResourceState srcState, - SubresourceRange srcSubresource, - ITextureResource::Offset3D srcOffset, - ITextureResource::Extents extent) -{ - assert(srcSubresource.mipLevelCount <= 1); -} - -void ResourceCommandEncoder::textureSubresourceBarrier( - ITextureResource* texture, - SubresourceRange subresourceRange, - ResourceState src, - ResourceState dst) -{ + auto encoder = m_commandBuffer->getMetalBlitCommandEncoder(); + encoder->resolveCounters( + static_cast<QueryPoolImpl*>(queryPool)->m_counterSampleBuffer.get(), + NS::Range(index, count), + static_cast<BufferResourceImpl*>(buffer)->m_buffer.get(), + offset); } void ResourceCommandEncoder::beginDebugEvent(const char* name, float rgbColor[3]) { + NS::SharedPtr<NS::String> string = MetalUtil::createString(name); + m_commandBuffer->m_commandBuffer->pushDebugGroup(string.get()); } void ResourceCommandEncoder::endDebugEvent() { + m_commandBuffer->m_commandBuffer->popDebugGroup(); } void RenderCommandEncoder::beginPass(IRenderPassLayout* renderPass, IFramebuffer* framebuffer) @@ -145,7 +221,7 @@ void RenderCommandEncoder::beginPass(IRenderPassLayout* renderPass, IFramebuffer if (rpd->depthAttachment() && false) { TextureResourceViewImpl* depthView = static_cast<TextureResourceViewImpl*>(fb->depthStencilView.get()); - rpd->depthAttachment()->setTexture(depthView->m_texture->m_texture); + rpd->depthAttachment()->setTexture(depthView->m_texture->m_texture.get()); } const int colorTargetCount = fb->renderTargetViews.getCount(); for (int i = 0; i < colorTargetCount; ++i) @@ -153,15 +229,7 @@ void RenderCommandEncoder::beginPass(IRenderPassLayout* renderPass, IFramebuffer TextureResourceViewImpl* texView = static_cast<TextureResourceViewImpl*>(fb->renderTargetViews[i].get()); MTL::Texture* tex = nullptr; assert(texView->m_texture); - if (texView->m_texture->m_isCurrentDrawable) - { - CA::MetalDrawable* drawable = static_cast<CA::MetalDrawable*>(fb->m_renderer->m_drawable); - tex = drawable->texture(); - } - else - { - tex = texView->m_texture->m_texture; - } + tex = texView->m_texture->m_texture.get(); rpd->colorAttachments()->object(i)->setTexture(tex); rpd->colorAttachments()->object(i)->setClearColor(MTL::ClearColor(0.2, 0.4, 0.9, 1.0)); } @@ -181,7 +249,7 @@ Result RenderCommandEncoder::bindPipeline( { m_currentPipeline = static_cast<PipelineStateImpl*>(pipelineState); // Initialize the root object - SLANG_RETURN_ON_FAIL(m_commandBuffer->m_rootObject.init(m_commandBuffer->m_renderer, + SLANG_RETURN_ON_FAIL(m_commandBuffer->m_rootObject.init(m_commandBuffer->m_device, m_currentPipeline->getProgram<ShaderProgramImpl>()->m_rootObjectLayout)); *outRootObject = &m_commandBuffer->m_rootObject; //if (pPipelineState->m_renderState == nullptr) return SLANG_ERROR_INVALID_PARAMETER; @@ -249,8 +317,8 @@ void RenderCommandEncoder::setVertexBuffers( BufferResourceImpl* buffer = static_cast<BufferResourceImpl*>(buffers[i]); if (buffer) { - MTL::Buffer* vertexBuffers = {buffer->m_buffer}; - m_encoder->setVertexBuffer(buffer->m_buffer, offsets[i], slotIndex); + MTL::Buffer* vertexBuffers = {buffer->m_buffer.get()}; + m_encoder->setVertexBuffer(buffer->m_buffer.get(), offsets[i], slotIndex); // ... } } @@ -369,12 +437,15 @@ Result RenderCommandEncoder::drawMeshTasks(int x, int y, int z) return SLANG_E_NOT_IMPLEMENTED; } -void ComputeCommandEncoder::endEncoding() { } +void ComputeCommandEncoder::endEncoding() +{ + ResourceCommandEncoder::endEncoding(); +} Result ComputeCommandEncoder::bindPipeline( IPipelineState* pipelineState, IShaderObject** outRootObject) { - return SLANG_E_NOT_IMPLEMENTED; + return setPipelineStateImpl(pipelineState, outRootObject); } Result ComputeCommandEncoder::bindPipelineWithRootObject( @@ -386,10 +457,15 @@ Result ComputeCommandEncoder::bindPipelineWithRootObject( Result ComputeCommandEncoder::dispatchCompute(int x, int y, int z) { auto pipeline = static_cast<PipelineStateImpl*>(m_currentPipeline.Ptr()); - if (!pipeline) - { - return SLANG_FAIL; - } + pipeline->ensureAPIPipelineStateCreated(); + + auto metalComputeCommandEncoder = m_commandBuffer->getMetalComputeCommandEncoder(); + metalComputeCommandEncoder->setComputePipelineState(pipeline->m_computePipelineState.get()); + ComputeBindingContext bindingContext; + bindingContext.init(m_commandBuffer->m_device, metalComputeCommandEncoder); + auto program = static_cast<ShaderProgramImpl*>(m_currentPipeline->m_program.get()); + m_commandBuffer->m_rootObject.bindAsRoot(&bindingContext, program->m_rootObjectLayout); + metalComputeCommandEncoder->dispatchThreadgroups(MTL::Size(x, y, z), pipeline->m_threadGroupSize); // Also create descriptor sets based on the given pipeline layout return SLANG_E_NOT_IMPLEMENTED; diff --git a/tools/gfx/metal/metal-command-encoder.h b/tools/gfx/metal/metal-command-encoder.h index 863018ca4..851ee60af 100644 --- a/tools/gfx/metal/metal-command-encoder.h +++ b/tools/gfx/metal/metal-command-encoder.h @@ -15,12 +15,15 @@ namespace metal class PipelineCommandEncoder : public ComObject { public: - - void init(CommandBufferImpl* commandBuffer); CommandBufferImpl* m_commandBuffer; - MTL::CommandBuffer* m_metalCommandBuffer = nullptr; - DeviceImpl* m_device = nullptr; + MTL::CommandBuffer* m_metalCommandBuffer; RefPtr<PipelineStateImpl> m_currentPipeline; + + void init(CommandBufferImpl* commandBuffer); + void endEncodingImpl(); + + Result setPipelineStateImpl(IPipelineState* state, IShaderObject** outRootObject); + }; class ResourceCommandEncoder @@ -47,28 +50,17 @@ public: virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } + virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; + + virtual SLANG_NO_THROW void SLANG_MCALL + writeTimestamp(IQueryPool* queryPool, GfxIndex index) override; + virtual SLANG_NO_THROW void SLANG_MCALL copyBuffer( IBufferResource* dst, Offset dstOffset, IBufferResource* src, Offset srcOffset, Size size) override; - virtual SLANG_NO_THROW void SLANG_MCALL - uploadBufferData(IBufferResource* buffer, Offset offset, Size size, void* data) override; - virtual SLANG_NO_THROW void SLANG_MCALL textureBarrier( - GfxCount count, - ITextureResource* const* textures, - ResourceState src, - ResourceState dst) override; - virtual SLANG_NO_THROW void SLANG_MCALL bufferBarrier( - GfxCount count, - IBufferResource* const* buffers, - ResourceState src, - ResourceState dst) override; - virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; - - virtual SLANG_NO_THROW void SLANG_MCALL - writeTimestamp(IQueryPool* queryPool, GfxIndex index) override; virtual SLANG_NO_THROW void SLANG_MCALL copyTexture( ITextureResource* dst, @@ -81,6 +73,20 @@ public: ITextureResource::Offset3D srcOffset, ITextureResource::Extents extent) override; + virtual SLANG_NO_THROW void SLANG_MCALL copyTextureToBuffer( + IBufferResource* dst, + Offset dstOffset, + Size dstSize, + Size dstRowStride, + ITextureResource* src, + ResourceState srcState, + SubresourceRange srcSubresource, + ITextureResource::Offset3D srcOffset, + ITextureResource::Extents extent) override; + + virtual SLANG_NO_THROW void SLANG_MCALL + uploadBufferData(IBufferResource* buffer, Offset offset, Size size, void* data) override; + virtual SLANG_NO_THROW void SLANG_MCALL uploadTextureData( ITextureResource* dst, SubresourceRange subResourceRange, @@ -89,6 +95,24 @@ public: ITextureResource::SubresourceData* subResourceData, GfxCount subResourceDataCount) override; + virtual SLANG_NO_THROW void SLANG_MCALL bufferBarrier( + GfxCount count, + IBufferResource* const* buffers, + ResourceState src, + ResourceState dst) override; + + virtual SLANG_NO_THROW void SLANG_MCALL textureBarrier( + GfxCount count, + ITextureResource* const* textures, + ResourceState src, + ResourceState dst) override; + + virtual SLANG_NO_THROW void SLANG_MCALL textureSubresourceBarrier( + ITextureResource* texture, + SubresourceRange subresourceRange, + ResourceState src, + ResourceState dst) override; + void _clearColorImage(TextureResourceViewImpl* viewImpl, ClearValue* clearValue); void _clearDepthImage( @@ -114,22 +138,7 @@ public: IBufferResource* buffer, Offset offset) override; - virtual SLANG_NO_THROW void SLANG_MCALL copyTextureToBuffer( - IBufferResource* dst, - Offset dstOffset, - Size dstSize, - Size dstRowStride, - ITextureResource* src, - ResourceState srcState, - SubresourceRange srcSubresource, - ITextureResource::Offset3D srcOffset, - ITextureResource::Extents extent) override; - virtual SLANG_NO_THROW void SLANG_MCALL textureSubresourceBarrier( - ITextureResource* texture, - SubresourceRange subresourceRange, - ResourceState src, - ResourceState dst) override; virtual SLANG_NO_THROW void SLANG_MCALL beginDebugEvent(const char* name, float rgbColor[3]) override; diff --git a/tools/gfx/metal/metal-command-queue.cpp b/tools/gfx/metal/metal-command-queue.cpp index c8b36ff1e..920b6ef7e 100644 --- a/tools/gfx/metal/metal-command-queue.cpp +++ b/tools/gfx/metal/metal-command-queue.cpp @@ -3,6 +3,7 @@ #include "metal-command-buffer.h" #include "metal-fence.h" +#include "metal-util.h" namespace gfx { @@ -23,22 +24,21 @@ CommandQueueImpl::~CommandQueueImpl() { } -void CommandQueueImpl::init(DeviceImpl* renderer) +void CommandQueueImpl::init(DeviceImpl* device, NS::SharedPtr<MTL::CommandQueue> commandQueue) { - m_renderer = renderer; - - MTL::Device* device = m_renderer->m_device; - m_commandQueue = device->newCommandQueue(8); + m_device = device; + m_commandQueue = commandQueue; } void CommandQueueImpl::waitOnHost() { + // TODO implement } Result CommandQueueImpl::getNativeHandle(InteropHandle* outHandle) { outHandle->api = InteropHandleAPI::Metal; - outHandle->handleValue = reinterpret_cast<intptr_t>(m_commandQueue); + outHandle->handleValue = reinterpret_cast<intptr_t>(m_commandQueue.get()); return SLANG_OK; } @@ -47,23 +47,57 @@ const CommandQueueImpl::Desc& CommandQueueImpl::getDesc() { return m_desc; } Result CommandQueueImpl::waitForFenceValuesOnDevice( GfxCount fenceCount, IFence** fences, uint64_t* waitValues) { - return SLANG_E_NOT_IMPLEMENTED; + for (GfxCount i = 0; i < fenceCount; ++i) + { + FenceWaitInfo waitInfo; + waitInfo.fence = static_cast<FenceImpl*>(fences[i]); + waitInfo.waitValue = waitValues[i]; + m_pendingWaitFences.add(waitInfo); + } + return SLANG_OK; } void CommandQueueImpl::queueSubmitImpl( uint32_t count, ICommandBuffer* const* commandBuffers, IFence* fence, uint64_t valueToSignal) { + // If there are any pending wait fences, encode them to a new command buffer. + // Metal ensures that command buffers are executed in the order they are committed. + if (m_pendingWaitFences.getCount() > 0) + { + MTL::CommandBuffer* commandBuffer = m_commandQueue->commandBuffer(); + for (const auto& fenceInfo : m_pendingWaitFences) + { + commandBuffer->encodeWait(fenceInfo.fence->m_event.get(), fenceInfo.waitValue); + } + commandBuffer->commit(); + m_pendingWaitFences.clear(); + } + for (uint32_t i = 0; i < count; ++i) { CommandBufferImpl* cmdBufImpl = static_cast<CommandBufferImpl*>(commandBuffers[i]); - cmdBufImpl->m_commandBuffer->presentDrawable(m_renderer->m_drawable); + // If this is the last command buffer and a fence is provided, signal the fence. + if (i == count - 1 && fence != nullptr) + { + cmdBufImpl->m_commandBuffer->encodeSignalEvent(static_cast<FenceImpl*>(fence)->m_event.get(), valueToSignal); + } cmdBufImpl->m_commandBuffer->commit(); } + + // If there are no command buffers to submit, but a fence is provided, signal the fence. + if (count == 0 && fence != nullptr) + { + MTL::CommandBuffer* commandBuffer = m_commandQueue->commandBuffer(); + commandBuffer->encodeSignalEvent(static_cast<FenceImpl*>(fence)->m_event.get(), valueToSignal); + commandBuffer->commit(); + } } void CommandQueueImpl::executeCommandBuffers( GfxCount count, ICommandBuffer* const* commandBuffers, IFence* fence, uint64_t valueToSignal) { + AUTORELEASEPOOL + if (count == 0 && fence == nullptr) { return; diff --git a/tools/gfx/metal/metal-command-queue.h b/tools/gfx/metal/metal-command-queue.h index 298b0a09a..045f4ed73 100644 --- a/tools/gfx/metal/metal-command-queue.h +++ b/tools/gfx/metal/metal-command-queue.h @@ -21,12 +21,20 @@ public: ICommandQueue* getInterface(const Guid& guid); public: + RefPtr<DeviceImpl> m_device; Desc m_desc; - RefPtr<DeviceImpl> m_renderer; - MTL::CommandQueue* m_commandQueue = nullptr; + NS::SharedPtr<MTL::CommandQueue> m_commandQueue; + + struct FenceWaitInfo + { + RefPtr<FenceImpl> fence; + uint64_t waitValue; + }; + List<FenceWaitInfo> m_pendingWaitFences; + ~CommandQueueImpl(); - void init(DeviceImpl* renderer); + void init(DeviceImpl* device, NS::SharedPtr<MTL::CommandQueue> commandQueue); virtual SLANG_NO_THROW void SLANG_MCALL waitOnHost() override; diff --git a/tools/gfx/metal/metal-device.cpp b/tools/gfx/metal/metal-device.cpp index f20ca5d1f..d844e6e89 100644 --- a/tools/gfx/metal/metal-device.cpp +++ b/tools/gfx/metal/metal-device.cpp @@ -10,14 +10,14 @@ #include "metal-shader-program.h" #include "metal-buffer.h" //#include "metal-command-queue.h" -//#include "metal-fence.h" -//#include "metal-query.h" +#include "metal-fence.h" +#include "metal-query.h" //#include "metal-resource-views.h" -//#include "metal-sampler.h" +#include "metal-sampler.h" #include "metal-shader-object.h" #include "metal-shader-object-layout.h" //#include "metal-shader-table.h" -//#include "metal-transient-heap.h" +#include "metal-transient-heap.h" //#include "metal-pipeline-dump-layer.h" //#include "metal-helper-functions.h" @@ -43,13 +43,15 @@ DeviceImpl::~DeviceImpl() Result DeviceImpl::getNativeDeviceHandles(InteropHandles* outHandles) { - outHandles->handles[0].handleValue = reinterpret_cast<intptr_t>(m_device); outHandles->handles[0].api = InteropHandleAPI::Metal; + outHandles->handles[0].handleValue = reinterpret_cast<intptr_t>(m_device.get()); return SLANG_OK; } SlangResult DeviceImpl::initialize(const Desc& desc) { + AUTORELEASEPOOL + // Initialize device info. { m_info.apiName = "Metal"; @@ -66,15 +68,15 @@ SlangResult DeviceImpl::initialize(const Desc& desc) SLANG_RETURN_ON_FAIL(RendererBase::initialize(desc)); SlangResult initDeviceResult = SLANG_OK; - m_device = MTL::CreateSystemDefaultDevice(); - m_commandQueue = m_device->newCommandQueue(); + m_device = NS::TransferPtr(MTL::CreateSystemDefaultDevice()); + m_commandQueue = NS::TransferPtr(m_device->newCommandQueue(64)); SLANG_RETURN_ON_FAIL(slangContext.initialize( desc.slang, desc.extendedDescCount, desc.extendedDescs, - SLANG_METAL, - "sm_5_1", + SLANG_METAL_LIB, + "", makeArray(slang::PreprocessorMacroDesc{ "__METAL__", "1" }).getView())); // TODO: expose via some other means @@ -89,11 +91,10 @@ SlangResult DeviceImpl::initialize(const Desc& desc) exit(1); } d->setDestination(MTL::CaptureDestinationGPUTraceDocument); - d->setCaptureObject(m_device); - std::string cpath("frame.gputrace"); - NS::String* path = NS::String::alloc()->init(cpath.c_str(), NS::UTF8StringEncoding); - NS::URL* url = NS::URL::alloc()->initFileURLWithPath(path); - d->setOutputURL(url); + d->setCaptureObject(m_device.get()); + NS::SharedPtr<NS::String> path = MetalUtil::createString("frame.gputrace"); + NS::SharedPtr<NS::URL> url = NS::TransferPtr(NS::URL::alloc()->initFileURLWithPath(path.get())); + d->setOutputURL(url.get()); NS::Error* errorCode = NS::Error::alloc(); if (!captureManager->startCapture(d, &errorCode)) { @@ -109,23 +110,30 @@ SlangResult DeviceImpl::initialize(const Desc& desc) //void DeviceImpl::waitForGpu() { m_deviceQueue.flushAndWait(); } -SLANG_NO_THROW const DeviceInfo& SLANG_MCALL DeviceImpl::getDeviceInfo() const { return m_info; } +const DeviceInfo& DeviceImpl::getDeviceInfo() const +{ + return m_info; +} -Result DeviceImpl::createTransientResourceHeap( - const ITransientResourceHeap::Desc& desc, ITransientResourceHeap** outHeap) +Result DeviceImpl::createTransientResourceHeap(const ITransientResourceHeap::Desc& desc, ITransientResourceHeap** outHeap) { + AUTORELEASEPOOL + RefPtr<TransientResourceHeapImpl> result = new TransientResourceHeapImpl(); - SLANG_RETURN_ON_FAIL(result->init(this, desc)); + SLANG_RETURN_ON_FAIL(result->init(desc, this)); returnComPtr(outHeap, result); return SLANG_OK; } Result DeviceImpl::createCommandQueue(const ICommandQueue::Desc& desc, ICommandQueue** outQueue) { + AUTORELEASEPOOL + if (m_queueAllocCount != 0) return SLANG_FAIL; + RefPtr<CommandQueueImpl> result = new CommandQueueImpl; - result->init(this); + result->init(this, m_commandQueue); returnComPtr(outQueue, result); m_queueAllocCount++; return SLANG_OK; @@ -134,24 +142,28 @@ Result DeviceImpl::createCommandQueue(const ICommandQueue::Desc& desc, ICommandQ Result DeviceImpl::createSwapchain( const ISwapchain::Desc& desc, WindowHandle window, ISwapchain** outSwapchain) { + AUTORELEASEPOOL + RefPtr<SwapchainImpl> sc = new SwapchainImpl(); SLANG_RETURN_ON_FAIL(sc->init(this, desc, window)); returnComPtr(outSwapchain, sc); return SLANG_OK; } -Result DeviceImpl::createFramebufferLayout( - const IFramebufferLayout::Desc& desc, IFramebufferLayout** outLayout) +Result DeviceImpl::createFramebufferLayout(const IFramebufferLayout::Desc& desc, IFramebufferLayout** outLayout) { + AUTORELEASEPOOL + RefPtr<FramebufferLayoutImpl> layout = new FramebufferLayoutImpl; SLANG_RETURN_ON_FAIL(layout->init(this, desc)); returnComPtr(outLayout, layout); return SLANG_OK; } -Result DeviceImpl::createRenderPassLayout( - const IRenderPassLayout::Desc& desc, IRenderPassLayout** outRenderPassLayout) +Result DeviceImpl::createRenderPassLayout(const IRenderPassLayout::Desc& desc, IRenderPassLayout** outRenderPassLayout) { + AUTORELEASEPOOL + RefPtr<RenderPassLayoutImpl> result = new RenderPassLayoutImpl; SLANG_RETURN_ON_FAIL(result->init(this, desc)); returnComPtr(outRenderPassLayout, result); @@ -160,6 +172,8 @@ Result DeviceImpl::createRenderPassLayout( Result DeviceImpl::createFramebuffer(const IFramebuffer::Desc& desc, IFramebuffer** outFramebuffer) { + AUTORELEASEPOOL + RefPtr<FramebufferImpl> fb = new FramebufferImpl; SLANG_RETURN_ON_FAIL(fb->init(this, desc)); returnComPtr(outFramebuffer, fb); @@ -173,36 +187,98 @@ SlangResult DeviceImpl::readTextureResource( Size* outRowPitch, Size* outPixelSize) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } SlangResult DeviceImpl::readBufferResource( - IBufferResource* inBuffer, Offset offset, Size size, ISlangBlob** outBlob) + IBufferResource* buffer, Offset offset, Size size, ISlangBlob** outBlob) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + // create staging buffer + NS::SharedPtr<MTL::Buffer> stagingBuffer = NS::TransferPtr(m_device->newBuffer(size, MTL::StorageModeShared)); + if (!stagingBuffer) + { + return SLANG_FAIL; + } + + MTL::CommandBuffer* commandBuffer = m_commandQueue->commandBuffer(); + MTL::BlitCommandEncoder* blitEncoder = commandBuffer->blitCommandEncoder(); + blitEncoder->copyFromBuffer(static_cast<BufferResourceImpl*>(buffer)->m_buffer.get(), offset, stagingBuffer.get(), 0, size); + blitEncoder->endEncoding(); + commandBuffer->commit(); + commandBuffer->waitUntilCompleted(); + + List<uint8_t> blobData; + blobData.setCount(size); + ::memcpy(blobData.getBuffer(), stagingBuffer->contents(), size); + auto blob = ListBlob::moveCreate(blobData); + + returnComPtr(outBlob, blob); + return SLANG_OK; } Result DeviceImpl::getAccelerationStructurePrebuildInfo( const IAccelerationStructure::BuildInputs& buildInputs, IAccelerationStructure::PrebuildInfo* outPrebuildInfo) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::createAccelerationStructure( const IAccelerationStructure::CreateDesc& desc, IAccelerationStructure** outAS) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::getTextureAllocationInfo( const ITextureResource::Desc& descIn, Size* outSize, Size* outAlignment) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + auto alignTo = [&](Size size, Size alignment) -> Size { + return ((size + alignment - 1) / alignment) * alignment; + }; + + TextureResource::Desc desc = fixupTextureDesc(descIn); + FormatInfo formatInfo; + gfxGetFormatInfo(desc.format, &formatInfo); + MTL::PixelFormat pixelFormat = MetalUtil::translatePixelFormat(desc.format); + Size alignment = m_device->minimumLinearTextureAlignmentForPixelFormat(pixelFormat); + Size size = 0; + ITextureResource::Extents extents = desc.size; + extents.width = extents.width ? extents.width : 1; + extents.height = extents.height ? extents.height : 1; + extents.depth = extents.depth ? extents.depth : 1; + + for (Int i = 0; i < desc.numMipLevels; ++i) + { + Size rowSize = ((extents.width + formatInfo.blockWidth - 1) / formatInfo.blockWidth) * formatInfo.blockSizeInBytes; + rowSize = alignTo(rowSize, alignment); + Size sliceSize = rowSize * alignTo(extents.height, formatInfo.blockHeight); + size += sliceSize * extents.depth; + extents.width = Math::Max(1, extents.width / 2); + extents.height = Math::Max(1, extents.height / 2); + extents.depth = Math::Max(1, extents.depth / 2); + } + size *= desc.arraySize ? desc.arraySize : 1; + + *outSize = size; + *outAlignment = alignment; + + return SLANG_OK; } Result DeviceImpl::getTextureRowAlignment(Size* outAlignment) { + AUTORELEASEPOOL + *outAlignment = 1; return SLANG_E_NOT_IMPLEMENTED; } @@ -212,107 +288,176 @@ Result DeviceImpl::createTextureResource( const ITextureResource::SubresourceData* initData, ITextureResource** outResource) { + AUTORELEASEPOOL + TextureResource::Desc desc = fixupTextureDesc(descIn); - const MTL::PixelFormat format = MetalUtil::getMetalPixelFormat(desc.format); - if (format == MTL::PixelFormat::PixelFormatInvalid) + const MTL::PixelFormat pixelFormat = MetalUtil::translatePixelFormat(desc.format); + if (pixelFormat == MTL::PixelFormat::PixelFormatInvalid) { assert(!"Unsupported texture format"); return SLANG_FAIL; } - RefPtr<TextureResourceImpl> textureResource(new TextureResourceImpl(desc, this)); - //textureResource->m_metalFormat = format; - MTL::TextureDescriptor* metalDesc = MTL::TextureDescriptor::alloc()->init(); - metalDesc->setStorageMode(MTL::StorageMode::StorageModePrivate); - // Create the texture + RefPtr<TextureResourceImpl> textureImpl(new TextureResourceImpl(desc, this)); + + NS::SharedPtr<MTL::TextureDescriptor> textureDesc = NS::TransferPtr(MTL::TextureDescriptor::alloc()->init()); + switch (desc.memoryType) + { + case MemoryType::DeviceLocal: + textureDesc->setStorageMode(MTL::StorageModePrivate); + break; + case MemoryType::Upload: + textureDesc->setStorageMode(MTL::StorageModeShared); + textureDesc->setCpuCacheMode(MTL::CPUCacheModeWriteCombined); + break; + case MemoryType::ReadBack: + textureDesc->setStorageMode(MTL::StorageModeShared); + break; + } + + NS::UInteger arrayLength = calcEffectiveArraySize(desc); + switch (desc.type) { case IResource::Type::Texture1D: - { - metalDesc->setTextureType(MTL::TextureType::TextureType1D); - metalDesc->setWidth(descIn.size.width); + textureDesc->setTextureType(arrayLength > 1 ? MTL::TextureType1DArray : MTL::TextureType1D); + textureDesc->setWidth(desc.size.width); break; - } case IResource::Type::Texture2D: - { - metalDesc->setTextureType(MTL::TextureType::TextureType2D); - metalDesc->setWidth(descIn.size.width); - metalDesc->setHeight(descIn.size.height); + if (desc.sampleDesc.numSamples > 1) + { + textureDesc->setTextureType(arrayLength > 1 ? MTL::TextureType2DMultisampleArray : MTL::TextureType2DMultisample); + textureDesc->setSampleCount(desc.sampleDesc.numSamples); + } + else + { + textureDesc->setTextureType(arrayLength > 1 ? MTL::TextureType2DArray : MTL::TextureType2D); + } + textureDesc->setWidth(descIn.size.width); + textureDesc->setHeight(descIn.size.height); break; - } case IResource::Type::TextureCube: - { - metalDesc->setTextureType(MTL::TextureType::TextureTypeCube); - metalDesc->setWidth(descIn.size.width); - metalDesc->setHeight(descIn.size.height); + textureDesc->setTextureType(arrayLength > 6 ? MTL::TextureTypeCubeArray : MTL::TextureTypeCube); + textureDesc->setWidth(descIn.size.width); + textureDesc->setHeight(descIn.size.height); break; - } case IResource::Type::Texture3D: - { - metalDesc->setTextureType(MTL::TextureType::TextureType3D); - metalDesc->setWidth(descIn.size.width); - metalDesc->setHeight(descIn.size.height); - metalDesc->setDepth(descIn.size.depth); + textureDesc->setTextureType(MTL::TextureType::TextureType3D); + textureDesc->setWidth(descIn.size.width); + textureDesc->setHeight(descIn.size.height); + textureDesc->setDepth(descIn.size.depth); break; - } default: - { assert("!Unsupported texture type"); return SLANG_FAIL; } + + MTL::TextureUsage textureUsage = MTL::TextureUsageUnknown; + if (desc.allowedStates.contains(ResourceState::RenderTarget)) + { + textureUsage |= MTL::TextureUsageRenderTarget; } - metalDesc->setMipmapLevelCount(desc.numMipLevels); - const int arraySize(calcEffectiveArraySize(desc)); - metalDesc->setArrayLength(arraySize); - metalDesc->setPixelFormat(format); - //metalDesc.setResourceOptions(); - metalDesc->setUsage(MTL::TextureUsageUnknown); - metalDesc->setSampleCount(desc.sampleDesc.numSamples); - textureResource->m_texture = m_device->newTexture(metalDesc); - - returnComPtr(outResource, textureResource); + if (desc.allowedStates.contains(ResourceState::ShaderResource)) + { + textureUsage |= MTL::TextureUsageShaderRead; + } + if (desc.allowedStates.contains(ResourceState::UnorderedAccess)) + { + textureUsage |= MTL::TextureUsageShaderWrite; + textureUsage |= MTL::TextureUsageShaderAtomic; + } + + textureDesc->setMipmapLevelCount(desc.numMipLevels); + textureDesc->setArrayLength(arrayLength); + textureDesc->setPixelFormat(pixelFormat); + textureDesc->setUsage(textureUsage); + textureDesc->setSampleCount(desc.sampleDesc.numSamples); + textureDesc->setAllowGPUOptimizedContents(desc.memoryType == MemoryType::DeviceLocal); + + textureImpl->m_texture = NS::TransferPtr(m_device->newTexture(textureDesc.get())); + if (!textureImpl->m_texture) + { + return SLANG_FAIL; + } + + returnComPtr(outResource, textureImpl); return SLANG_OK; } Result DeviceImpl::createBufferResource( const IBufferResource::Desc& descIn, const void* initData, IBufferResource** outResource) { + AUTORELEASEPOOL + BufferResource::Desc desc = fixupBufferDesc(descIn); const Size bufferSize = desc.sizeInBytes; - MTL::ResourceOptions opts = (desc.memoryType == MemoryType::DeviceLocal ? MTL::ResourceStorageModePrivate : 0); - - RefPtr<BufferResourceImpl> buffer(new BufferResourceImpl(desc, this)); + MTL::ResourceOptions resourceOptions = MTL::ResourceOptions(0); + switch (desc.memoryType) + { + case MemoryType::DeviceLocal: + resourceOptions = MTL::ResourceStorageModePrivate; + break; + case MemoryType::Upload: + resourceOptions = MTL::ResourceStorageModeShared | MTL::CPUCacheModeWriteCombined; + break; + case MemoryType::ReadBack: + resourceOptions = MTL::ResourceStorageModeShared; + break; + } + resourceOptions |= (desc.memoryType == MemoryType::DeviceLocal) ? MTL::ResourceStorageModePrivate : MTL::ResourceStorageModeShared; - if (initData) + RefPtr<BufferResourceImpl> bufferImpl(new BufferResourceImpl(desc, this)); + bufferImpl->m_buffer = NS::TransferPtr(m_device->newBuffer(bufferSize, resourceOptions)); + if (!bufferImpl->m_buffer) { - buffer->m_buffer = m_device->newBuffer(initData, bufferSize, opts); + return SLANG_FAIL; } - else + + if (initData) { - buffer->m_buffer = m_device->newBuffer(bufferSize, opts); + NS::SharedPtr<MTL::Buffer> stagingBuffer = NS::TransferPtr(m_device->newBuffer( + initData, bufferSize, MTL::ResourceStorageModeShared | MTL::CPUCacheModeWriteCombined)); + MTL::CommandBuffer* commandBuffer = m_commandQueue->commandBuffer(); + MTL::BlitCommandEncoder* encoder = commandBuffer->blitCommandEncoder(); + if (!stagingBuffer || !commandBuffer || !encoder) + { + return SLANG_FAIL; + } + encoder->copyFromBuffer(stagingBuffer.get(), 0, bufferImpl->m_buffer.get(), 0, bufferSize); + encoder->endEncoding(); + commandBuffer->commit(); } - returnComPtr(outResource, buffer); + returnComPtr(outResource, bufferImpl); return SLANG_OK; } Result DeviceImpl::createBufferFromNativeHandle( InteropHandle handle, const IBufferResource::Desc& srcDesc, IBufferResource** outResource) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::createSamplerState(ISamplerState::Desc const& desc, ISamplerState** outSampler) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<SamplerStateImpl> samplerImpl = new SamplerStateImpl(); + SLANG_RETURN_ON_FAIL(samplerImpl->init(this, desc)); + returnComPtr(outSampler, samplerImpl); + return SLANG_OK; } Result DeviceImpl::createTextureView( ITextureResource* texture, IResourceView::Desc const& desc, IResourceView** outView) { + AUTORELEASEPOOL + auto resourceImpl = static_cast<TextureResourceImpl*>(texture); RefPtr<TextureResourceViewImpl> view = new TextureResourceViewImpl(this); view->m_desc = desc; @@ -325,7 +470,7 @@ Result DeviceImpl::createTextureView( } bool isArray = resourceImpl->getDesc()->arraySize > 1; - MTL::PixelFormat pixelFormat = MetalUtil::getMetalPixelFormat(desc.format); + MTL::PixelFormat pixelFormat = MetalUtil::translatePixelFormat(desc.format); NS::Range levelRange(desc.subresourceRange.baseArrayLayer, std::max(desc.subresourceRange.layerCount, 1)); NS::Range sliceRange(desc.subresourceRange.mipLevel, std::max(desc.subresourceRange.mipLevelCount, 1)); MTL::TextureType textureType; @@ -356,8 +501,7 @@ Result DeviceImpl::createTextureView( view->m_type = ResourceViewImpl::ViewType::Texture; view->m_texture = resourceImpl; //new TextureResourceImpl(newDesc, this); - view->m_texture->m_isCurrentDrawable = resourceImpl->m_isCurrentDrawable; - view->m_texture->m_texture = resourceImpl->m_texture->newTextureView(pixelFormat, textureType, levelRange, sliceRange); + view->m_texture->m_texture = NS::TransferPtr(resourceImpl->m_texture->newTextureView(pixelFormat, textureType, levelRange, sliceRange)); returnComPtr(outView, view); return SLANG_OK; @@ -365,7 +509,29 @@ Result DeviceImpl::createTextureView( Result DeviceImpl::getFormatSupportedResourceStates(Format format, ResourceStateSet* outStates) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + // TODO - add table based on https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf + ResourceStateSet allowedStates; + allowedStates.add(ResourceState::VertexBuffer); + allowedStates.add(ResourceState::IndexBuffer); + allowedStates.add(ResourceState::ConstantBuffer); + allowedStates.add(ResourceState::ShaderResource); + allowedStates.add(ResourceState::UnorderedAccess); + allowedStates.add(ResourceState::RenderTarget); + allowedStates.add(ResourceState::DepthRead); + allowedStates.add(ResourceState::DepthWrite); + allowedStates.add(ResourceState::Present); + allowedStates.add(ResourceState::IndirectArgument); + allowedStates.add(ResourceState::CopySource); + allowedStates.add(ResourceState::ResolveSource); + allowedStates.add(ResourceState::CopyDestination); + allowedStates.add(ResourceState::ResolveDestination); + allowedStates.add(ResourceState::AccelerationStructure); + allowedStates.add(ResourceState::AccelerationStructureBuildInput); + + *outStates = allowedStates; + return SLANG_OK; } Result DeviceImpl::createBufferView( @@ -374,76 +540,88 @@ Result DeviceImpl::createBufferView( IResourceView::Desc const& desc, IResourceView** outView) { - return SLANG_E_NOT_IMPLEMENTED; -} + AUTORELEASEPOOL -static MTL::VertexStepFunction translateVertexStepFunction(const InputSlotClass& slotClass) -{ - switch (slotClass) + // Counter buffers are not supported on metal. + if (counterBuffer) { - case InputSlotClass::PerInstance: return MTL::VertexStepFunctionPerInstance; - case InputSlotClass::PerVertex: - default: return MTL::VertexStepFunctionPerVertex; + return SLANG_FAIL; } -} -Result DeviceImpl::createInputLayout(IInputLayout::Desc const& desc, IInputLayout** outLayout) -{ - RefPtr<InputLayoutImpl> layout(new InputLayoutImpl); - List<MTL::VertexDescriptor*>& dstAttributes = layout->m_vertexDescs; - List<MTL::VertexBufferLayoutDescriptor*>& dstBufferLayouts = layout->m_bufferLayoutDescs; + if (desc.type != IResourceView::Type::UnorderedAccess && desc.type != IResourceView::Type::ShaderResource) + { + return SLANG_FAIL; + } - const InputElementDesc* srcElements = desc.inputElements; - Int numElements = desc.inputElementCount; + auto bufferImpl = static_cast<BufferResourceImpl*>(buffer); - const VertexStreamDesc* srcVertexStreams = desc.vertexStreams; - Int vertexStreamCount = desc.vertexStreamCount; + RefPtr<BufferResourceViewImpl> viewImpl = new BufferResourceViewImpl(this); + viewImpl->m_desc = desc; + viewImpl->m_buffer = bufferImpl; + viewImpl->m_offset = desc.bufferRange.offset;; + viewImpl->m_size = desc.bufferRange.size == 0 ? bufferImpl->getDesc()->sizeInBytes : desc.bufferRange.size;; + returnComPtr(outView, viewImpl); + return SLANG_OK; +} - dstAttributes.setCount(numElements); - dstBufferLayouts.setCount(vertexStreamCount); +Result DeviceImpl::createInputLayout(IInputLayout::Desc const& desc, IInputLayout** outLayout) +{ + AUTORELEASEPOOL - for (Int i = 0; i < vertexStreamCount; ++i) + RefPtr<InputLayoutImpl> layout(new InputLayoutImpl); + layout->m_vertexDescriptor = NS::TransferPtr(MTL::VertexDescriptor::alloc()->init()); + if (!layout->m_vertexDescriptor) { - auto& vbld = dstBufferLayouts[i]; - auto& srcStream = srcVertexStreams[i]; - vbld->setStepFunction(translateVertexStepFunction(srcStream.slotClass)); - vbld->setStepRate(srcStream.instanceDataStepRate); - vbld->setStride(srcStream.stride); + return SLANG_FAIL; } - for (Int i = 0; i < numElements; ++i) + for (Int i = 0; i < desc.inputElementCount; ++i) { - auto& srcAttrib = srcElements[i]; - auto& dstAttrib = dstAttributes[i]; - dstAttrib->attributes()->object(i)->setOffset(srcAttrib.offset); - dstAttrib->attributes()->object(i)->setBufferIndex(srcAttrib.bufferSlotIndex); - MTL::VertexFormat metalFormat = MetalUtil::getMetalVertexFormat(srcAttrib.format); + const auto& inputElement = desc.inputElements[i]; + MTL::VertexAttributeDescriptor* desc = layout->m_vertexDescriptor->attributes()->object(i); + desc->setOffset(inputElement.offset); + desc->setBufferIndex(inputElement.bufferSlotIndex); + MTL::VertexFormat metalFormat = MetalUtil::translateVertexFormat(inputElement.format); if (metalFormat == MTL::VertexFormatInvalid) { return SLANG_FAIL; } - dstAttrib->attributes()->object(i)->setFormat(metalFormat); + desc->setFormat(metalFormat); } + for (Int i = 0; i < desc.vertexStreamCount; ++i) + { + const auto& vertexStream = desc.vertexStreams[i]; + MTL::VertexBufferLayoutDescriptor* desc = layout->m_vertexDescriptor->layouts()->object(i); + desc->setStepFunction(MetalUtil::translateVertexStepFunction(vertexStream.slotClass)); + desc->setStepRate(vertexStream.instanceDataStepRate); + desc->setStride(vertexStream.stride); + } + + returnComPtr(outLayout, layout); return SLANG_OK; } Result DeviceImpl::createProgram( const IShaderProgram::Desc& desc, IShaderProgram** outProgram, ISlangBlob** outDiagnosticBlob) { - // TODO: + AUTORELEASEPOOL + RefPtr<ShaderProgramImpl> shaderProgram = new ShaderProgramImpl(this); shaderProgram->init(desc); - //m_deviceObjectsWithPotentialBackReferences.add(shaderProgram); - - RootShaderObjectLayout::create( + RootShaderObjectLayoutImpl::create( this, shaderProgram->linkedProgram, shaderProgram->linkedProgram->getLayout(), shaderProgram->m_rootObjectLayout.writeRef()); - returnComPtr(outProgram, shaderProgram); + if (!shaderProgram->isSpecializable()) + { + SLANG_RETURN_ON_FAIL(shaderProgram->compileShaders(this)); + } + + returnComPtr(outProgram, shaderProgram); return SLANG_OK; } @@ -452,62 +630,94 @@ Result DeviceImpl::createShaderObjectLayout( slang::TypeLayoutReflection* typeLayout, ShaderObjectLayoutBase** outLayout) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<ShaderObjectLayoutImpl> layout; + SLANG_RETURN_ON_FAIL(ShaderObjectLayoutImpl::createForElementType( + this, session, typeLayout, layout.writeRef())); + returnRefPtrMove(outLayout, layout); + return SLANG_OK; } Result DeviceImpl::createShaderObject(ShaderObjectLayoutBase* layout, IShaderObject** outObject) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<ShaderObjectImpl> shaderObject; + SLANG_RETURN_ON_FAIL(ShaderObjectImpl::create(this, + static_cast<ShaderObjectLayoutImpl*>(layout), shaderObject.writeRef())); + returnComPtr(outObject, shaderObject); + return SLANG_OK; } Result DeviceImpl::createMutableShaderObject( ShaderObjectLayoutBase* layout, IShaderObject** outObject) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::createMutableRootShaderObject(IShaderProgram* program, IShaderObject** outObject) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::createShaderTable(const IShaderTable::Desc& desc, IShaderTable** outShaderTable) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } -Result DeviceImpl::createGraphicsPipelineState( - const GraphicsPipelineStateDesc& inDesc, IPipelineState** outState) +Result DeviceImpl::createGraphicsPipelineState(const GraphicsPipelineStateDesc& desc, IPipelineState** outState) { - GraphicsPipelineStateDesc desc = inDesc; + AUTORELEASEPOOL + RefPtr<PipelineStateImpl> pipelineStateImpl = new PipelineStateImpl(this); pipelineStateImpl->init(desc); - pipelineStateImpl->establishStrongDeviceReference(); - //m_deviceObjectsWithPotentialBackReferences.add(pipelineStateImpl); returnComPtr(outState, pipelineStateImpl); return SLANG_OK; } -Result DeviceImpl::createComputePipelineState( - const ComputePipelineStateDesc& inDesc, IPipelineState** outState) +Result DeviceImpl::createComputePipelineState(const ComputePipelineStateDesc& desc, IPipelineState** outState) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<PipelineStateImpl> pipelineStateImpl = new PipelineStateImpl(this); + pipelineStateImpl->init(desc); + m_deviceObjectsWithPotentialBackReferences.add(pipelineStateImpl); + returnComPtr(outState, pipelineStateImpl); + return SLANG_OK; } -Result DeviceImpl::createRayTracingPipelineState( - const RayTracingPipelineStateDesc& desc, IPipelineState** outState) +Result DeviceImpl::createRayTracingPipelineState(const RayTracingPipelineStateDesc& desc, IPipelineState** outState) { + AUTORELEASEPOOL + return SLANG_E_NOT_IMPLEMENTED; } Result DeviceImpl::createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outPool) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<QueryPoolImpl> poolImpl = new QueryPoolImpl(); + SLANG_RETURN_ON_FAIL(poolImpl->init(this, desc)); + returnComPtr(outPool, poolImpl); + return SLANG_OK; } Result DeviceImpl::createFence(const IFence::Desc& desc, IFence** outFence) { - return SLANG_E_NOT_IMPLEMENTED; + AUTORELEASEPOOL + + RefPtr<FenceImpl> fenceImpl = new FenceImpl(); + SLANG_RETURN_ON_FAIL(fenceImpl->init(this, desc)); + returnComPtr(outFence, fenceImpl); + return SLANG_OK; } Result DeviceImpl::waitForFences( diff --git a/tools/gfx/metal/metal-device.h b/tools/gfx/metal/metal-device.h index 4e536ddc4..6c3ba68f8 100644 --- a/tools/gfx/metal/metal-device.h +++ b/tools/gfx/metal/metal-device.h @@ -3,7 +3,6 @@ #include "metal-base.h" #include "metal-device.h" -#include "metal-command-buffer.h" #include "../simple-transient-resource-heap.h" #include "metal-framebuffer.h" @@ -19,7 +18,6 @@ class DeviceImpl : public RendererBase { public: // Renderer implementation - using TransientResourceHeapImpl = SimpleTransientResourceHeap<DeviceImpl, CommandBufferImpl>; virtual SLANG_NO_THROW Result SLANG_MCALL initialize(const Desc& desc) override; virtual SLANG_NO_THROW Result SLANG_MCALL getFormatSupportedResourceStates(Format format, ResourceStateSet* outStates) override; @@ -128,18 +126,14 @@ public: ~DeviceImpl(); public: - DeviceInfo m_info; String m_adapterName; - MTL::CaptureManager* m_captureManager = nullptr; - MTL::Drawable* m_drawable = nullptr; - CA::MetalLayer* m_metalLayer = nullptr; bool captureEnabled() const { return std::getenv("MTL_CAPTURE") != nullptr; } Desc m_desc; - MTL::Device* m_device = nullptr; - MTL::CommandQueue* m_commandQueue = nullptr; + NS::SharedPtr<MTL::Device> m_device; + NS::SharedPtr<MTL::CommandQueue> m_commandQueue; //DescriptorSetAllocator descriptorSetAllocator; diff --git a/tools/gfx/metal/metal-fence.cpp b/tools/gfx/metal/metal-fence.cpp index 3c0e8edef..c8c318d7f 100644 --- a/tools/gfx/metal/metal-fence.cpp +++ b/tools/gfx/metal/metal-fence.cpp @@ -10,38 +10,44 @@ using namespace Slang; namespace metal { -FenceImpl::FenceImpl(DeviceImpl* device) - : m_device(device) -{} - FenceImpl::~FenceImpl() { } -Result FenceImpl::init(const IFence::Desc& desc) +Result FenceImpl::init(DeviceImpl* device, const IFence::Desc& desc) { - return SLANG_FAIL; + m_device = device; + m_event = NS::TransferPtr(m_device->m_device->newSharedEvent()); + if (!m_event) + { + return SLANG_FAIL; + } + m_event->setSignaledValue(desc.initialValue); + return SLANG_OK; } Result FenceImpl::getCurrentValue(uint64_t* outValue) { - return SLANG_E_NOT_IMPLEMENTED; + *outValue = m_event->signaledValue(); + return SLANG_OK; } Result FenceImpl::setCurrentValue(uint64_t value) { - return SLANG_E_NOT_IMPLEMENTED; + m_event->setSignaledValue(value); + return SLANG_OK; } Result FenceImpl::getSharedHandle(InteropHandle* outHandle) { - return SLANG_E_NOT_IMPLEMENTED; + return SLANG_E_NOT_AVAILABLE; } Result FenceImpl::getNativeHandle(InteropHandle* outNativeHandle) { - outNativeHandle->handleValue = 0; - return SLANG_FAIL; + outNativeHandle->api = InteropHandleAPI::Metal; + outNativeHandle->handleValue = reinterpret_cast<intptr_t>(m_event.get()); + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-fence.h b/tools/gfx/metal/metal-fence.h index af5cb7806..d8aed7dfe 100644 --- a/tools/gfx/metal/metal-fence.h +++ b/tools/gfx/metal/metal-fence.h @@ -15,21 +15,19 @@ class FenceImpl : public FenceBase { public: RefPtr<DeviceImpl> m_device; - - FenceImpl(DeviceImpl* device); + NS::SharedPtr<MTL::SharedEvent> m_event; ~FenceImpl(); - Result init(const IFence::Desc& desc); + Result init(DeviceImpl* device, const IFence::Desc& desc); virtual SLANG_NO_THROW Result SLANG_MCALL getCurrentValue(uint64_t* outValue) override; virtual SLANG_NO_THROW Result SLANG_MCALL setCurrentValue(uint64_t value) override; - + virtual SLANG_NO_THROW Result SLANG_MCALL getSharedHandle(InteropHandle* outHandle) override; - virtual SLANG_NO_THROW Result SLANG_MCALL - getNativeHandle(InteropHandle* outNativeHandle) override; + virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outNativeHandle) override; }; } // namespace metal diff --git a/tools/gfx/metal/metal-framebuffer.cpp b/tools/gfx/metal/metal-framebuffer.cpp index c676f44eb..46e60b2eb 100644 --- a/tools/gfx/metal/metal-framebuffer.cpp +++ b/tools/gfx/metal/metal-framebuffer.cpp @@ -17,11 +17,11 @@ FramebufferLayoutImpl::~FramebufferLayoutImpl() //m_renderPass->release(); } -Result FramebufferLayoutImpl::init(DeviceImpl* renderer, const IFramebufferLayout::Desc& desc) +Result FramebufferLayoutImpl::init(DeviceImpl* device, const IFramebufferLayout::Desc& desc) { // Metal doesn't have a notion of Framebuffers or FramebufferLayouts per se. // We simply stash the desc and use it when creating the (convenience) Framebuffer - m_renderer = renderer; + m_device = device; m_desc = desc; return SLANG_OK; } @@ -30,9 +30,9 @@ FramebufferImpl::~FramebufferImpl() { } -Result FramebufferImpl::init(DeviceImpl* renderer, const IFramebuffer::Desc& desc) +Result FramebufferImpl::init(DeviceImpl* device, const IFramebuffer::Desc& desc) { - m_renderer = renderer; + m_device = device; m_layout = static_cast<FramebufferLayoutImpl*>(desc.layout); m_width = m_height = 1; diff --git a/tools/gfx/metal/metal-framebuffer.h b/tools/gfx/metal/metal-framebuffer.h index 10ee637b8..7bd3b0bef 100644 --- a/tools/gfx/metal/metal-framebuffer.h +++ b/tools/gfx/metal/metal-framebuffer.h @@ -20,6 +20,8 @@ enum class FramebufferLayoutImpl : public FramebufferLayoutBase { public: + RefPtr<DeviceImpl> m_device; + Desc m_desc; #if 0 MTL::RenderPassDescriptor* m_renderPass = nullptr; Array<MTL::RenderPassColorAttachmentDescriptor*, kMaxTargets> m_targetDescs; @@ -28,22 +30,20 @@ public: bool m_hasDepthStencilTarget = false; uint32_t m_renderTargetCount = 0; #endif - DeviceImpl* m_renderer = nullptr; - Desc m_desc; public: ~FramebufferLayoutImpl(); - Result init(DeviceImpl* renderer, const IFramebufferLayout::Desc& desc); + Result init(DeviceImpl* device, const IFramebufferLayout::Desc& desc); }; class FramebufferImpl : public FramebufferBase { public: + BreakableReference<DeviceImpl> m_device; ShortList<ComPtr<IResourceView>> renderTargetViews; ComPtr<IResourceView> depthStencilView; uint32_t m_width; uint32_t m_height; - BreakableReference<DeviceImpl> m_renderer; MTL::ClearColor m_clearValues[kMaxTargets]; RefPtr<FramebufferLayoutImpl> m_layout; #if 0 @@ -55,7 +55,7 @@ public: public: ~FramebufferImpl(); - Result init(DeviceImpl* renderer, const IFramebuffer::Desc& desc); + Result init(DeviceImpl* device, const IFramebuffer::Desc& desc); }; } // namespace metal diff --git a/tools/gfx/metal/metal-helper-functions.cpp b/tools/gfx/metal/metal-helper-functions.cpp index e884f2e76..69c4327ad 100644 --- a/tools/gfx/metal/metal-helper-functions.cpp +++ b/tools/gfx/metal/metal-helper-functions.cpp @@ -1,12 +1,45 @@ // metal-helper-functions.cpp #include "metal-helper-functions.h" #include "metal-device.h" +#include "metal-util.h" namespace gfx { using namespace Slang; +Result SLANG_MCALL getMetalAdapters(List<AdapterInfo>& outAdapters) +{ + AUTORELEASEPOOL + + auto addAdapter = [&](MTL::Device* device) + { + AdapterInfo info = {}; + const char* name = device->name()->cString(NS::ASCIIStringEncoding); + memcpy(info.name, name, Math::Min(strlen(name), sizeof(AdapterInfo::name) - 1)); + uint64_t registryID = device->registryID(); + memcpy(&info.luid.luid[0], ®istryID, sizeof(registryID)); + outAdapters.add(info); + }; + + NS::Array* devices = MTL::CopyAllDevices(); + if (devices->count() > 0) + { + for (int i = 0; i < devices->count(); ++i) + { + MTL::Device* device = static_cast<MTL::Device*>(devices->object(i)); + addAdapter(device); + } + } + else + { + MTL::Device* device = MTL::CreateSystemDefaultDevice(); + addAdapter(device); + device->release(); + } + return SLANG_OK; +} + Result SLANG_MCALL createMetalDevice(const IDevice::Desc* desc, IDevice** outRenderer) { RefPtr<metal::DeviceImpl> result = new metal::DeviceImpl(); diff --git a/tools/gfx/metal/metal-helper-functions.h b/tools/gfx/metal/metal-helper-functions.h index 94b005321..0936cfa46 100644 --- a/tools/gfx/metal/metal-helper-functions.h +++ b/tools/gfx/metal/metal-helper-functions.h @@ -4,10 +4,96 @@ namespace gfx { - using namespace Slang; +namespace metal +{ + +/// A "simple" binding offset that records an offset in buffer/texture/sampler slots +struct BindingOffset +{ + uint32_t buffer = 0; + uint32_t texture = 0; + uint32_t sampler = 0; + + /// Create a default (zero) offset + BindingOffset() = default; + + /// Create an offset based on offset information in the given Slang `varLayout` + BindingOffset(slang::VariableLayoutReflection* varLayout) + { + if (varLayout) + { + buffer = (uint32_t)varLayout->getOffset(SLANG_PARAMETER_CATEGORY_METAL_BUFFER); + texture = (uint32_t)varLayout->getOffset(SLANG_PARAMETER_CATEGORY_METAL_TEXTURE); + sampler = (uint32_t)varLayout->getOffset(SLANG_PARAMETER_CATEGORY_METAL_SAMPLER); + } + } + + /// Create an offset based on size/stride information in the given Slang `typeLayout` + BindingOffset(slang::TypeLayoutReflection* typeLayout) + { + if (typeLayout) + { + buffer = (uint32_t)typeLayout->getSize(SLANG_PARAMETER_CATEGORY_METAL_BUFFER); + texture = (uint32_t)typeLayout->getSize(SLANG_PARAMETER_CATEGORY_METAL_TEXTURE); + sampler = (uint32_t)typeLayout->getSize(SLANG_PARAMETER_CATEGORY_METAL_SAMPLER); + } + } + + /// Add any values in the given `offset` + void operator+=(BindingOffset const& offset) + { + buffer += offset.buffer; + texture += offset.texture; + sampler += offset.sampler; + } +}; + +/// Contextual data and operations required when binding shader objects to the pipeline state +struct BindingContext +{ + DeviceImpl* device = nullptr; + virtual void setData(const void* data, NS::UInteger length, NS::UInteger index) = 0; + virtual void setBuffer(MTL::Buffer* buffer, NS::UInteger index) = 0; + virtual void setTexture(MTL::Texture* texture, NS::UInteger index) = 0; + virtual void setSampler(MTL::SamplerState* sampler, NS::UInteger index) = 0; +}; + +struct ComputeBindingContext : public BindingContext +{ + MTL::ComputeCommandEncoder* encoder; + + Result init(DeviceImpl* device, MTL::ComputeCommandEncoder* encoder) + { + this->device = device; + this->encoder = encoder; + return SLANG_OK; + } + + void setData(const void* data, NS::UInteger length, NS::UInteger index) override + { + encoder->setBytes(data, length, index); + } + + void setBuffer(MTL::Buffer* buffer, NS::UInteger index) override + { + encoder->setBuffer(buffer, 0, index); + } + + void setTexture(MTL::Texture* texture, NS::UInteger index) override + { + encoder->setTexture(texture, index); + } + + void setSampler(MTL::SamplerState* sampler, NS::UInteger index) override + { + encoder->setSamplerState(sampler, index); + } +}; + +} // namespace metal -//Result SLANG_MCALL getMetalAdapters(List<AdapterInfo>& outAdapters); +Result SLANG_MCALL getMetalAdapters(List<AdapterInfo>& outAdapters); Result SLANG_MCALL createMetalDevice(const IDevice::Desc* desc, IDevice** outRenderer); } // namespace gfx diff --git a/tools/gfx/metal/metal-pipeline-state.cpp b/tools/gfx/metal/metal-pipeline-state.cpp index aca10c12a..b0fce682e 100644 --- a/tools/gfx/metal/metal-pipeline-state.cpp +++ b/tools/gfx/metal/metal-pipeline-state.cpp @@ -5,6 +5,7 @@ #include "metal-shader-program.h" #include "metal-shader-object-layout.h" #include "metal-vertex-layout.h" +#include "metal-util.h" namespace gfx { @@ -15,51 +16,41 @@ namespace metal { PipelineStateImpl::PipelineStateImpl(DeviceImpl* device) + : m_device(device) { - // Only weakly reference `device` at start. - // We make it a strong reference only when the pipeline state is exposed to the user. - // Note that `PipelineState`s may also be created via implicit specialization that - // happens behind the scenes, and the user will not have access to those specialized - // pipeline states. Only those pipeline states that are returned to the user needs to - // hold a strong reference to `device`. - m_device.setWeakReference(device); } PipelineStateImpl::~PipelineStateImpl() { } -void PipelineStateImpl::establishStrongDeviceReference() { m_device.establishStrongReference(); } - -void PipelineStateImpl::comFree() { m_device.breakStrongReference(); } - -void PipelineStateImpl::init(const GraphicsPipelineStateDesc& inDesc) +void PipelineStateImpl::init(const GraphicsPipelineStateDesc& desc) { PipelineStateDesc pipelineDesc; pipelineDesc.type = PipelineType::Graphics; - pipelineDesc.graphics = inDesc; + pipelineDesc.graphics = desc; initializeBase(pipelineDesc); } -void PipelineStateImpl::init(const ComputePipelineStateDesc& inDesc) +void PipelineStateImpl::init(const ComputePipelineStateDesc& desc) { PipelineStateDesc pipelineDesc; pipelineDesc.type = PipelineType::Compute; - pipelineDesc.compute = inDesc; + pipelineDesc.compute = desc; initializeBase(pipelineDesc); } -void PipelineStateImpl::init(const RayTracingPipelineStateDesc& inDesc) +void PipelineStateImpl::init(const RayTracingPipelineStateDesc& desc) { PipelineStateDesc pipelineDesc; pipelineDesc.type = PipelineType::RayTracing; - pipelineDesc.rayTracing.set(inDesc); + pipelineDesc.rayTracing.set(desc); initializeBase(pipelineDesc); } Result PipelineStateImpl::createMetalRenderPipelineState() { - MTL::RenderPipelineDescriptor* pd = MTL::RenderPipelineDescriptor::alloc()->init(); + NS::SharedPtr<MTL::RenderPipelineDescriptor> pd = NS::TransferPtr(MTL::RenderPipelineDescriptor::alloc()->init()); auto programImpl = static_cast<ShaderProgramImpl*>(m_program.Ptr()); if (programImpl) { @@ -94,8 +85,8 @@ Result PipelineStateImpl::createMetalRenderPipelineState() // Set default rasterization state // Set default framebuffer layout NS::Error* error; - m_renderState = m_device->m_device->newRenderPipelineState(pd, &error); - if (m_renderState == nullptr) + m_renderPipelineState = NS::TransferPtr(m_device->m_device->newRenderPipelineState(pd.get(), &error)); + if (!m_renderPipelineState) { std::cout << error->localizedDescription()->utf8String() << std::endl; return SLANG_E_INVALID_ARG; @@ -105,20 +96,36 @@ Result PipelineStateImpl::createMetalRenderPipelineState() Result PipelineStateImpl::createMetalComputePipelineState() { - return SLANG_E_NOT_IMPLEMENTED; + auto programImpl = static_cast<ShaderProgramImpl*>(m_program.Ptr()); + if (!programImpl) + return SLANG_FAIL; + + NS::SharedPtr<MTL::ComputePipelineDescriptor> pd = NS::TransferPtr(MTL::ComputePipelineDescriptor::alloc()->init()); + + auto functionName = MetalUtil::createString(programImpl->m_entryPointNames[0].getBuffer()); + NS::SharedPtr<MTL::Function> function = NS::TransferPtr(programImpl->m_modules[0]->newFunction(functionName.get())); + if (!function) + return SLANG_FAIL; + + NS::Error *error; + m_computePipelineState = NS::TransferPtr(m_device->m_device->newComputePipelineState(function.get(), &error)); + + // Query thread group size for use during dispatch. + SlangUInt threadGroupSize[3]; + programImpl->linkedProgram->getLayout()->getEntryPointByIndex(0)->getComputeThreadGroupSize(3, threadGroupSize); + m_threadGroupSize = MTL::Size(threadGroupSize[0], threadGroupSize[1], threadGroupSize[2]); + + return m_computePipelineState ? SLANG_OK : SLANG_FAIL; } Result PipelineStateImpl::ensureAPIPipelineStateCreated() { - if (m_renderState) - return SLANG_OK; - switch (desc.type) { case PipelineType::Compute: - return createMetalComputePipelineState(); + return m_computePipelineState ? SLANG_OK : createMetalComputePipelineState(); case PipelineType::Graphics: - return createMetalRenderPipelineState(); + return m_renderPipelineState ? SLANG_OK : createMetalRenderPipelineState(); default: SLANG_UNREACHABLE("Unknown pipeline type."); return SLANG_FAIL; diff --git a/tools/gfx/metal/metal-pipeline-state.h b/tools/gfx/metal/metal-pipeline-state.h index 321c110a1..e816af5b1 100644 --- a/tools/gfx/metal/metal-pipeline-state.h +++ b/tools/gfx/metal/metal-pipeline-state.h @@ -14,19 +14,17 @@ namespace metal class PipelineStateImpl : public PipelineStateBase { public: + RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::RenderPipelineState> m_renderPipelineState; + NS::SharedPtr<MTL::ComputePipelineState> m_computePipelineState; + MTL::Size m_threadGroupSize; + PipelineStateImpl(DeviceImpl* device); ~PipelineStateImpl(); - // Turns `m_device` into a strong reference. - // This method should be called before returning the pipeline state object to - // external users (i.e. via an `IPipelineState` pointer). - void establishStrongDeviceReference(); - - virtual void comFree() override; - - void init(const GraphicsPipelineStateDesc& inDesc); - void init(const ComputePipelineStateDesc& inDesc); - void init(const RayTracingPipelineStateDesc& inDesc); + void init(const GraphicsPipelineStateDesc& desc); + void init(const ComputePipelineStateDesc& desc); + void init(const RayTracingPipelineStateDesc& desc); Result createMetalComputePipelineState(); Result createMetalRenderPipelineState(); @@ -34,11 +32,6 @@ public: virtual Result ensureAPIPipelineStateCreated() override; virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override; - - BreakableReference<DeviceImpl> m_device; - - MTL::RenderPipelineState* m_renderState = nullptr; - MTL::ComputePipelineState* m_computeState = nullptr; }; class RayTracingPipelineStateImpl : public PipelineStateImpl diff --git a/tools/gfx/metal/metal-query.cpp b/tools/gfx/metal/metal-query.cpp index bfed117c6..4e1d09d77 100644 --- a/tools/gfx/metal/metal-query.cpp +++ b/tools/gfx/metal/metal-query.cpp @@ -10,13 +10,57 @@ using namespace Slang; namespace metal { -Result QueryPoolImpl::init(const IQueryPool::Desc& desc, DeviceImpl* device) + +QueryPoolImpl::~QueryPoolImpl() { - return SLANG_E_NOT_IMPLEMENTED; } -QueryPoolImpl::~QueryPoolImpl() +static MTL::CounterSet* findCounterSet(MTL::Device* device, QueryType queryType) +{ + if (queryType != QueryType::Timestamp) + { + return nullptr; + } + + static NS::String* timestampStr = MTLSTR("timestamp"); + + for (int i = 0; i < device->counterSets()->count(); ++i) + { + MTL::CounterSet* counterSet = static_cast<MTL::CounterSet*>(device->counterSets()->object(i)); + for (int j = 0; j < counterSet->counters()->count(); ++j) + { + MTL::Counter* counter = static_cast<MTL::Counter*>(counterSet->counters()->object(j)); + if (counter->name()->isEqualToString(MTL::CommonCounterTimestamp)) + { + return counterSet; + } + } + } + return nullptr; +} + +Result QueryPoolImpl::init(DeviceImpl* device, const IQueryPool::Desc& desc) { + m_device = device; + m_desc = desc; + + MTL::CounterSet* counterSet = findCounterSet(m_device->m_device.get(), m_desc.type); + if (!counterSet) + { + return SLANG_E_NOT_AVAILABLE; + } + + NS::SharedPtr<MTL::CounterSampleBufferDescriptor> counterSampleBufferDesc = NS::TransferPtr(MTL::CounterSampleBufferDescriptor::alloc()->init()); + counterSampleBufferDesc->setStorageMode(MTL::StorageModeShared); + counterSampleBufferDesc->setSampleCount(m_desc.count); + counterSampleBufferDesc->setCounterSet(counterSet); + + m_device->m_device->counterSets(); + + NS::Error* error; + m_counterSampleBuffer = NS::TransferPtr(m_device->m_device->newCounterSampleBuffer(counterSampleBufferDesc.get(), &error)); + + return m_counterSampleBuffer ? SLANG_OK : SLANG_FAIL; } Result QueryPoolImpl::getResult(GfxIndex index, GfxCount count, uint64_t* data) diff --git a/tools/gfx/metal/metal-query.h b/tools/gfx/metal/metal-query.h index 02c8d3172..3b3e28489 100644 --- a/tools/gfx/metal/metal-query.h +++ b/tools/gfx/metal/metal-query.h @@ -15,15 +15,16 @@ namespace metal class QueryPoolImpl : public QueryPoolBase { public: - Result init(const IQueryPool::Desc& desc, DeviceImpl* device); + RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::CounterSampleBuffer> m_counterSampleBuffer; + ~QueryPoolImpl(); -public: + Result init(DeviceImpl* device, const IQueryPool::Desc& desc); + virtual SLANG_NO_THROW Result SLANG_MCALL getResult(GfxIndex index, GfxCount count, uint64_t* data) override; -public: - RefPtr<DeviceImpl> m_device; }; } // namespace metal diff --git a/tools/gfx/metal/metal-render-pass.cpp b/tools/gfx/metal/metal-render-pass.cpp index b7016ebc2..f0849eb4b 100644 --- a/tools/gfx/metal/metal-render-pass.cpp +++ b/tools/gfx/metal/metal-render-pass.cpp @@ -48,15 +48,15 @@ static inline MTL::StoreAction translateStoreOp(IRenderPassLayout::TargetStoreOp } } -Result RenderPassLayoutImpl::init(DeviceImpl* renderer, const IRenderPassLayout::Desc& desc) +Result RenderPassLayoutImpl::init(DeviceImpl* device, const IRenderPassLayout::Desc& desc) { - m_renderer = renderer; + m_device = device; FramebufferLayoutImpl* framebufferLayout = static_cast<FramebufferLayoutImpl*>(desc.framebufferLayout); assert(framebufferLayout); // Initialize render pass descriptor, filling in attachment metadata, but leaving texture data unbound. - m_renderPassDesc = MTL::RenderPassDescriptor::alloc()->init(); + m_renderPassDesc = NS::TransferPtr(MTL::RenderPassDescriptor::alloc()->init()); m_renderPassDesc->setRenderTargetArrayLength(desc.renderTargetCount); MTL::RenderPassColorAttachmentDescriptorArray* colorAttachments = m_renderPassDesc->colorAttachments(); diff --git a/tools/gfx/metal/metal-render-pass.h b/tools/gfx/metal/metal-render-pass.h index 0ccf3424b..ef69d8909 100644 --- a/tools/gfx/metal/metal-render-pass.h +++ b/tools/gfx/metal/metal-render-pass.h @@ -21,11 +21,12 @@ public: IRenderPassLayout* getInterface(const Guid& guid); public: - MTL::RenderPassDescriptor* m_renderPassDesc = nullptr; - RefPtr<DeviceImpl> m_renderer; + RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::RenderPassDescriptor> m_renderPassDesc; + ~RenderPassLayoutImpl(); - Result init(DeviceImpl* renderer, const IRenderPassLayout::Desc& desc); + Result init(DeviceImpl* device, const IRenderPassLayout::Desc& desc); }; } // namespace metal diff --git a/tools/gfx/metal/metal-resource-views.cpp b/tools/gfx/metal/metal-resource-views.cpp index 8b05b66a1..76c0f21ed 100644 --- a/tools/gfx/metal/metal-resource-views.cpp +++ b/tools/gfx/metal/metal-resource-views.cpp @@ -18,6 +18,15 @@ Result TextureResourceViewImpl::getNativeHandle(InteropHandle* outHandle) return SLANG_E_NOT_IMPLEMENTED; } +BufferResourceViewImpl::~BufferResourceViewImpl() +{ +} + +Result BufferResourceViewImpl::getNativeHandle(InteropHandle* outHandle) +{ + return SLANG_E_NOT_IMPLEMENTED; +} + TexelBufferResourceViewImpl::TexelBufferResourceViewImpl(DeviceImpl* device) : ResourceViewImpl(ViewType::TexelBuffer, device) {} @@ -31,15 +40,6 @@ Result TexelBufferResourceViewImpl::getNativeHandle(InteropHandle* outHandle) return SLANG_E_NOT_IMPLEMENTED; } -PlainBufferResourceViewImpl::PlainBufferResourceViewImpl(DeviceImpl* device) - : ResourceViewImpl(ViewType::PlainBuffer, device) -{} - -Result PlainBufferResourceViewImpl::getNativeHandle(InteropHandle* outHandle) -{ - return m_buffer->getNativeResourceHandle(outHandle); -} - DeviceAddress AccelerationStructureImpl::getDeviceAddress() { return 0; diff --git a/tools/gfx/metal/metal-resource-views.h b/tools/gfx/metal/metal-resource-views.h index 735668ad1..9a8e24ab5 100644 --- a/tools/gfx/metal/metal-resource-views.h +++ b/tools/gfx/metal/metal-resource-views.h @@ -20,8 +20,8 @@ public: enum class ViewType { Texture, + Buffer, TexelBuffer, - PlainBuffer, }; public: @@ -45,21 +45,26 @@ public: virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override; }; -class TexelBufferResourceViewImpl : public ResourceViewImpl +class BufferResourceViewImpl : public ResourceViewImpl { public: - TexelBufferResourceViewImpl(DeviceImpl* device); - ~TexelBufferResourceViewImpl(); + BufferResourceViewImpl(DeviceImpl* device) + : ResourceViewImpl(ViewType::Buffer, device) + {} + ~BufferResourceViewImpl(); RefPtr<BufferResourceImpl> m_buffer; + Offset m_offset; + Size m_size; + virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override; }; -class PlainBufferResourceViewImpl : public ResourceViewImpl +class TexelBufferResourceViewImpl : public ResourceViewImpl { public: - PlainBufferResourceViewImpl(DeviceImpl* device); + TexelBufferResourceViewImpl(DeviceImpl* device); + ~TexelBufferResourceViewImpl(); RefPtr<BufferResourceImpl> m_buffer; - virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override; }; diff --git a/tools/gfx/metal/metal-sampler.cpp b/tools/gfx/metal/metal-sampler.cpp index cd94def37..561b93ea2 100644 --- a/tools/gfx/metal/metal-sampler.cpp +++ b/tools/gfx/metal/metal-sampler.cpp @@ -1,5 +1,6 @@ // metal-sampler.cpp #include "metal-sampler.h" +#include "metal-util.h" namespace gfx { @@ -9,17 +10,50 @@ using namespace Slang; namespace metal { -SamplerStateImpl::SamplerStateImpl(DeviceImpl* device) - : m_device(device) -{} - SamplerStateImpl::~SamplerStateImpl() { } +Result SamplerStateImpl::init(DeviceImpl* device, const ISamplerState::Desc& desc) +{ + m_device = device; + + NS::SharedPtr<MTL::SamplerDescriptor> samplerDesc = NS::TransferPtr(MTL::SamplerDescriptor::alloc()->init()); + + samplerDesc->setMinFilter(MetalUtil::translateSamplerMinMagFilter(desc.minFilter)); + samplerDesc->setMagFilter(MetalUtil::translateSamplerMinMagFilter(desc.magFilter)); + samplerDesc->setMipFilter(MetalUtil::translateSamplerMipFilter(desc.mipFilter)); + + samplerDesc->setSAddressMode(MetalUtil::translateSamplerAddressMode(desc.addressU)); + samplerDesc->setTAddressMode(MetalUtil::translateSamplerAddressMode(desc.addressV)); + samplerDesc->setRAddressMode(MetalUtil::translateSamplerAddressMode(desc.addressW)); + + samplerDesc->setMaxAnisotropy(Math::Clamp(desc.maxAnisotropy, 1u, 16u)); + + // TODO: support translation of border color... + MTL::SamplerBorderColor borderColor = MTL::SamplerBorderColorOpaqueBlack; + samplerDesc->setBorderColor(borderColor); + + samplerDesc->setNormalizedCoordinates(true); + + samplerDesc->setCompareFunction(MetalUtil::translateCompareFunction(desc.comparisonFunc)); + samplerDesc->setLodMinClamp(Math::Clamp(desc.minLOD, 0.f, 1000.f)); + samplerDesc->setLodMaxClamp(Math::Clamp(desc.maxLOD, samplerDesc->lodMinClamp(), 1000.f)); + + samplerDesc->setSupportArgumentBuffers(true); + + // TODO: no support for reduction op + + m_samplerState = NS::TransferPtr(m_device->m_device->newSamplerState(samplerDesc.get())); + + return m_samplerState ? SLANG_OK : SLANG_FAIL; +} + Result SamplerStateImpl::getNativeHandle(InteropHandle* outHandle) { - return SLANG_E_NOT_IMPLEMENTED; + outHandle->api = InteropHandleAPI::Metal; + outHandle->handleValue = reinterpret_cast<intptr_t>(m_samplerState.get()); + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-sampler.h b/tools/gfx/metal/metal-sampler.h index ac6b00e62..c7156cb2e 100644 --- a/tools/gfx/metal/metal-sampler.h +++ b/tools/gfx/metal/metal-sampler.h @@ -16,8 +16,12 @@ class SamplerStateImpl : public SamplerStateBase { public: RefPtr<DeviceImpl> m_device; - SamplerStateImpl(DeviceImpl* device); + NS::SharedPtr<MTL::SamplerState> m_samplerState; + ~SamplerStateImpl(); + + Result init(DeviceImpl* device, const ISamplerState::Desc& desc); + virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override; }; diff --git a/tools/gfx/metal/metal-shader-object-layout.cpp b/tools/gfx/metal/metal-shader-object-layout.cpp index 8439aa1ff..abde03b54 100644 --- a/tools/gfx/metal/metal-shader-object-layout.cpp +++ b/tools/gfx/metal/metal-shader-object-layout.cpp @@ -8,13 +8,310 @@ using namespace Slang; namespace metal { -Result RootShaderObjectLayout::create( - DeviceImpl* renderer, + +ShaderObjectLayoutImpl::SubObjectRangeOffset::SubObjectRangeOffset( + slang::VariableLayoutReflection* varLayout) + : BindingOffset(varLayout) +{ + if (auto pendingLayout = varLayout->getPendingDataLayout()) + { + pendingOrdinaryData = (uint32_t)pendingLayout->getOffset(SLANG_PARAMETER_CATEGORY_UNIFORM); + } +} + +ShaderObjectLayoutImpl::SubObjectRangeStride::SubObjectRangeStride( + slang::TypeLayoutReflection* typeLayout) + : BindingOffset(typeLayout) +{ + if (auto pendingLayout = typeLayout->getPendingDataTypeLayout()) + { + pendingOrdinaryData = (uint32_t)typeLayout->getStride(); + } +} + +Result ShaderObjectLayoutImpl::Builder::setElementTypeLayout(slang::TypeLayoutReflection* typeLayout) +{ + typeLayout = _unwrapParameterGroups(typeLayout, m_containerType); + + m_elementTypeLayout = typeLayout; + + m_totalOrdinaryDataSize = (uint32_t)typeLayout->getSize(); + if (m_totalOrdinaryDataSize > 0) + { + m_bufferCount++; + } + + // Compute the binding ranges that are used to store + // the logical contents of the object in memory. + + SlangInt bindingRangeCount = typeLayout->getBindingRangeCount(); + for (SlangInt r = 0; r < bindingRangeCount; ++r) + { + slang::BindingType slangBindingType = typeLayout->getBindingRangeType(r); + SlangInt count = typeLayout->getBindingRangeBindingCount(r); + slang::TypeLayoutReflection* slangLeafTypeLayout = + typeLayout->getBindingRangeLeafTypeLayout(r); + + BindingRangeInfo bindingRangeInfo; + bindingRangeInfo.bindingType = slangBindingType; + bindingRangeInfo.count = count; + switch (slangBindingType) + { + case slang::BindingType::ConstantBuffer: + case slang::BindingType::ParameterBlock: + case slang::BindingType::ExistentialValue: + bindingRangeInfo.baseIndex = m_subObjectCount; + bindingRangeInfo.subObjectIndex = m_subObjectCount; + m_subObjectCount += count; + break; + case slang::BindingType::RawBuffer: + case slang::BindingType::MutableRawBuffer: + bindingRangeInfo.baseIndex = m_bufferCount; + if (slangLeafTypeLayout->getType()->getElementType() != nullptr) + { + // A structured buffer occupies both a resource slot and + // a sub-object slot. + bindingRangeInfo.subObjectIndex = m_subObjectCount; + m_subObjectCount += count; + } + m_bufferCount += count; + m_bufferRanges.add(r); + break; + case slang::BindingType::Sampler: + bindingRangeInfo.baseIndex = m_samplerCount; + m_samplerCount += count; + m_samplerRanges.add(r); + break; + case slang::BindingType::Texture: + case slang::BindingType::MutableTexture: + bindingRangeInfo.baseIndex = m_textureCount; + m_textureCount += count; + m_textureRanges.add(r); + break; + case slang::BindingType::TypedBuffer: + case slang::BindingType::MutableTypedBuffer: + bindingRangeInfo.baseIndex = m_textureCount; + m_textureCount += count; + m_textureRanges.add(r); + break; + default: + break; + } + + // We'd like to extract the information on the Metal resource + // index that this range should bind into. + // + // A binding range represents a logical member of the shader + // object type, and it may encompass zero or more *descriptor + // ranges* that describe how it is physically bound to pipeline + // state. + // + // If the current binding range is backed by at least one descriptor + // range then we can query the register offset of that descriptor + // range. We expect that in the common case there will be exactly + // one descriptor range, and we can extract the information easily. + // + // TODO: we might eventually need to special-case our handling + // of combined texture-sampler ranges since they will need to + // store two different offsets. + // + if (typeLayout->getBindingRangeDescriptorRangeCount(r) != 0) + { + // The Slang reflection information organizes the descriptor ranges + // into "descriptor sets" but Metal has no notion like that so we + // expect all ranges belong to a single set. + // + SlangInt descriptorSetIndex = typeLayout->getBindingRangeDescriptorSetIndex(r); + SLANG_ASSERT(descriptorSetIndex == 0); + + SlangInt descriptorRangeIndex = typeLayout->getBindingRangeFirstDescriptorRangeIndex(r); + auto registerOffset = typeLayout->getDescriptorSetDescriptorRangeIndexOffset(descriptorSetIndex, descriptorRangeIndex); + + bindingRangeInfo.registerOffset = (uint32_t)registerOffset; + } + + m_bindingRanges.add(bindingRangeInfo); + } + + SlangInt subObjectRangeCount = typeLayout->getSubObjectRangeCount(); + for (SlangInt r = 0; r < subObjectRangeCount; ++r) + { + SlangInt bindingRangeIndex = typeLayout->getSubObjectRangeBindingRangeIndex(r); + auto& bindingRange = m_bindingRanges[bindingRangeIndex]; + + auto slangBindingType = typeLayout->getBindingRangeType(bindingRangeIndex); + slang::TypeLayoutReflection* slangLeafTypeLayout = + typeLayout->getBindingRangeLeafTypeLayout(bindingRangeIndex); + + SubObjectRangeInfo subObjectRange; + subObjectRange.bindingRangeIndex = bindingRangeIndex; + + // We will use Slang reflection information to extract the offset and stride + // information for each sub-object range. + // + subObjectRange.offset = SubObjectRangeOffset(typeLayout->getSubObjectRangeOffset(r)); + subObjectRange.stride = SubObjectRangeStride(slangLeafTypeLayout); + + // A sub-object range can either represent a sub-object of a known + // type, like a `ConstantBuffer<Foo>` or `ParameterBlock<Foo>` + // *or* it can represent a sub-object of some existential type (e.g., `IBar`). + // + RefPtr<ShaderObjectLayoutImpl> subObjectLayout; + switch (slangBindingType) + { + default: + { + // In the case of `ConstantBuffer<X>` or `ParameterBlock<X>` + // we can construct a layout from the element type directly. + // + auto elementTypeLayout = slangLeafTypeLayout->getElementTypeLayout(); + createForElementType( + m_renderer, + m_session, + elementTypeLayout, + subObjectLayout.writeRef()); + } + break; + case slang::BindingType::ExistentialValue: + // In the case of an interface-type sub-object range, we can only + // construct a layout if we have static specialization information + // that tells us what type we expect to find in that range. + // + // The static specialization information is expected to take the + // form of a "pending" type layotu attached to the interface type + // of the leaf type layout. + // + if (auto pendingTypeLayout = slangLeafTypeLayout->getPendingDataTypeLayout()) + { + createForElementType( + m_renderer, + m_session, + pendingTypeLayout, + subObjectLayout.writeRef()); + + // An interface-type range that includes ordinary data can + // increase the size of the ordinary data buffer we need to + // allocate for the parent object. + // + uint32_t ordinaryDataEnd = subObjectRange.offset.pendingOrdinaryData + + (uint32_t)bindingRange.count * subObjectRange.stride.pendingOrdinaryData; + + if (ordinaryDataEnd > m_totalOrdinaryDataSize) + { + m_totalOrdinaryDataSize = ordinaryDataEnd; + } + } + } + subObjectRange.layout = subObjectLayout; + + m_subObjectRanges.add(subObjectRange); + } + return SLANG_OK; +} + +SlangResult ShaderObjectLayoutImpl::Builder::build(ShaderObjectLayoutImpl** outLayout) +{ + auto layout = + RefPtr<ShaderObjectLayoutImpl>(new ShaderObjectLayoutImpl()); + SLANG_RETURN_ON_FAIL(layout->_init(this)); + + returnRefPtrMove(outLayout, layout); + return SLANG_OK; +} + +Result ShaderObjectLayoutImpl::createForElementType( + RendererBase* renderer, + slang::ISession* session, + slang::TypeLayoutReflection* elementType, + ShaderObjectLayoutImpl** outLayout) +{ + Builder builder(renderer, session); + builder.setElementTypeLayout(elementType); + return builder.build(outLayout); +} + +Result ShaderObjectLayoutImpl::_init(Builder const* builder) +{ + auto renderer = builder->m_renderer; + + initBase(renderer, builder->m_session, builder->m_elementTypeLayout); + + m_bindingRanges = builder->m_bindingRanges; + m_bufferRanges = builder->m_bufferRanges; + m_textureRanges = builder->m_textureRanges; + m_samplerRanges = builder->m_samplerRanges; + + m_bufferCount = builder->m_bufferCount; + m_textureCount = builder->m_textureCount; + m_samplerCount = builder->m_samplerCount; + m_subObjectCount = builder->m_subObjectCount; + m_subObjectRanges = builder->m_subObjectRanges; + + m_totalOrdinaryDataSize = builder->m_totalOrdinaryDataSize; + + m_containerType = builder->m_containerType; + return SLANG_OK; +} + +Result RootShaderObjectLayoutImpl::Builder::build(RootShaderObjectLayoutImpl** outLayout) +{ + RefPtr<RootShaderObjectLayoutImpl> layout = new RootShaderObjectLayoutImpl(); + SLANG_RETURN_ON_FAIL(layout->_init(this)); + + returnRefPtrMove(outLayout, layout); + return SLANG_OK; +} + +void RootShaderObjectLayoutImpl::Builder::addGlobalParams(slang::VariableLayoutReflection* globalsLayout) +{ + setElementTypeLayout(globalsLayout->getTypeLayout()); +} + +void RootShaderObjectLayoutImpl::Builder::addEntryPoint( + SlangStage stage, ShaderObjectLayoutImpl* entryPointLayout, slang::EntryPointLayout* slangEntryPoint) +{ + EntryPointInfo info; + info.layout = entryPointLayout; + info.offset = BindingOffset(slangEntryPoint->getVarLayout()); + m_entryPoints.add(info); +} + +Result RootShaderObjectLayoutImpl::create( + RendererBase* renderer, slang::IComponentType* program, slang::ProgramLayout* programLayout, - RootShaderObjectLayout** outLayout) + RootShaderObjectLayoutImpl** outLayout) { - return SLANG_E_NOT_IMPLEMENTED; + RootShaderObjectLayoutImpl::Builder builder(renderer, program, programLayout); + builder.addGlobalParams(programLayout->getGlobalParamsVarLayout()); + + SlangInt entryPointCount = programLayout->getEntryPointCount(); + for (SlangInt e = 0; e < entryPointCount; ++e) + { + auto slangEntryPoint = programLayout->getEntryPointByIndex(e); + RefPtr<ShaderObjectLayoutImpl> entryPointLayout; + SLANG_RETURN_ON_FAIL(ShaderObjectLayoutImpl::createForElementType( + renderer, program->getSession(), slangEntryPoint->getTypeLayout(), entryPointLayout.writeRef())); + builder.addEntryPoint(slangEntryPoint->getStage(), entryPointLayout, slangEntryPoint); + } + + SLANG_RETURN_ON_FAIL(builder.build(outLayout)); + + return SLANG_OK; +} + +Result RootShaderObjectLayoutImpl::_init(Builder const* builder) +{ + auto renderer = builder->m_renderer; + + SLANG_RETURN_ON_FAIL(Super::_init(builder)); + + m_program = builder->m_program; + m_programLayout = builder->m_programLayout; + m_entryPoints = builder->m_entryPoints; + m_slangSession = m_program->getSession(); + + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-shader-object-layout.h b/tools/gfx/metal/metal-shader-object-layout.h index 9d441f624..969c78c3a 100644 --- a/tools/gfx/metal/metal-shader-object-layout.h +++ b/tools/gfx/metal/metal-shader-object-layout.h @@ -2,7 +2,7 @@ #pragma once #include "metal-base.h" -#include "metal-device.h" +#include "metal-helper-functions.h" namespace gfx { @@ -12,36 +12,62 @@ using namespace Slang; namespace metal { -enum -{ - kMaxDescriptorSets = 32, -}; - class ShaderObjectLayoutImpl : public ShaderObjectLayoutBase { public: - struct BindingOffset - { - }; + // A shader object comprises three main kinds of state: + // + // * Zero or more bytes of ordinary ("uniform") data + // * Zero or more *bindings* for textures, buffers, and samplers + // * Zero or more *sub-objects* representing nested parameter blocks, etc. + // + // A shader object *layout* stores information that can be used to + // organize these different kinds of state and optimize access to them. + // + // For example, both texture/buffer/sampler bindings and sub-objects + // are organized into logical *binding ranges* by the Slang reflection + // API, and a shader object layout will store information about those + // ranges in a form that is usable for the Metal API: + + /// Information about a logical binding range as reported by Slang reflection struct BindingRangeInfo { + /// The type of bindings in this range slang::BindingType bindingType; + + /// The number of bindings in this range Index count; + + /// The starting index for this range in the appropriate "flat" array in a shader object. + /// E.g., for a buffers range, this would be an index into the `m_buffers` array. Index baseIndex; + + /// The offset of this binding range from the start of the sub-object. + uint32_t registerOffset; + + /// An index into the sub-object array if this binding range is treated + /// as a sub-object. Index subObjectIndex; - uint32_t bindingOffset; - uint32_t offset; + + /// TODO remove this once specialization is removed bool isSpecializable = false; }; + // Sometimes we just want to iterate over the ranges that represent + // sub-objects while skipping over the others, because sub-object + // ranges often require extra handling or more state. + // + // For that reason we also store pre-computed information about each + // sub-object range. + /// Offset information for a sub-object range struct SubObjectRangeOffset : BindingOffset { - SubObjectRangeOffset() {} + SubObjectRangeOffset() + {} + + SubObjectRangeOffset(slang::VariableLayoutReflection* varLayout); - SubObjectRangeOffset(slang::VariableLayoutReflection* varLayout) - // : BindingOffset(varLayout) - { } /// The offset for "pending" ordinary data related to this range uint32_t pendingOrdinaryData = 0; }; @@ -49,15 +75,10 @@ public: /// Stride information for a sub-object range struct SubObjectRangeStride : BindingOffset { - SubObjectRangeStride() {} + SubObjectRangeStride() + {} - SubObjectRangeStride(slang::TypeLayoutReflection* typeLayout) - { - if (auto pendingLayout = typeLayout->getPendingDataTypeLayout()) - { - pendingOrdinaryData = (uint32_t)pendingLayout->getStride(); - } - } + SubObjectRangeStride(slang::TypeLayoutReflection* typeLayout); /// The stride for "pending" ordinary data related to this range uint32_t pendingOrdinaryData = 0; @@ -79,41 +100,151 @@ public: SubObjectRangeStride stride; }; + struct Builder + { + public: + Builder(RendererBase* renderer, slang::ISession* session) + : m_renderer(renderer), m_session(session) + {} + + RendererBase* m_renderer; + slang::ISession* m_session; + slang::TypeLayoutReflection* m_elementTypeLayout; + + List<BindingRangeInfo> m_bindingRanges; + List<SubObjectRangeInfo> m_subObjectRanges; + + /// The indices of the binding ranges that represent buffers + List<Index> m_bufferRanges; + + /// The indices of the binding ranges that represent textures + List<Index> m_textureRanges; + + /// The indices of the binding ranges that represent samplers + List<Index> m_samplerRanges; + + Index m_bufferCount = 0; + Index m_textureCount = 0; + Index m_samplerCount = 0; + Index m_subObjectCount = 0; + + uint32_t m_totalOrdinaryDataSize = 0; + + /// The container type of this shader object. When `m_containerType` is + /// `StructuredBuffer` or `Array`, this shader object represents a collection + /// instead of a single object. + ShaderObjectContainerType m_containerType = ShaderObjectContainerType::None; + + Result setElementTypeLayout(slang::TypeLayoutReflection* typeLayout); + SlangResult build(ShaderObjectLayoutImpl** outLayout); + }; + + static Result createForElementType( + RendererBase* renderer, + slang::ISession* session, + slang::TypeLayoutReflection* elementType, + ShaderObjectLayoutImpl** outLayout); + + List<BindingRangeInfo> const& getBindingRanges() { return m_bindingRanges; } + + Index getBindingRangeCount() { return m_bindingRanges.getCount(); } + + BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRanges[index]; } - Index getBindingRangeCount() { return 0; } - BindingRangeInfo const& getBindingRange(Index index) { return m_bindingRangeInfo[index]; } - Index getResourceViewCount() { return 0; } - Index getSamplerCount() { return 0; } - Index getCombinedTextureSamplerCount() { return 0; } - Index getSubObjectCount() { return 0; } + Index getBufferCount() { return m_bufferCount; } + Index getTextureCount() { return m_textureCount; } + Index getSamplerCount() { return m_samplerCount; } + Index getSubObjectCount() { return m_subObjectCount; } SubObjectRangeInfo const& getSubObjectRange(Index index) { return m_subObjectRanges[index]; } List<SubObjectRangeInfo> const& getSubObjectRanges() { return m_subObjectRanges; } + + RendererBase* getRenderer() { return m_renderer; } + + slang::TypeReflection* getType() + { + return m_elementTypeLayout->getType(); + } + + /// Get the indices that represent all the buffer ranges in this type + List<Index> const& getBufferRanges() const { return m_bufferRanges; } + + /// Get the indices that reprsent all the texture ranges in this type + List<Index> const& getTextureRanges() const { return m_textureRanges; } + + /// Get the indices that represnet all the sampler ranges in this type + List<Index> const& getSamplerRanges() const { return m_samplerRanges; } + + uint32_t getTotalOrdinaryDataSize() const { return m_totalOrdinaryDataSize; } + protected: - List<BindingRangeInfo> m_bindingRangeInfo; - List<SubObjectRangeInfo> m_subObjectRanges; + Result _init(Builder const* builder); + + List<BindingRangeInfo> m_bindingRanges; + List<Index> m_bufferRanges; + List<Index> m_textureRanges; + List<Index> m_samplerRanges; + Index m_bufferCount = 0; + Index m_textureCount = 0; + Index m_samplerCount = 0; + Index m_subObjectCount = 0; + uint32_t m_totalOrdinaryDataSize = 0; + List<SubObjectRangeInfo> m_subObjectRanges; }; -class EntryPointLayout : public ShaderObjectLayoutImpl +class RootShaderObjectLayoutImpl : public ShaderObjectLayoutImpl { typedef ShaderObjectLayoutImpl Super; public: -}; + struct EntryPointInfo + { + RefPtr<ShaderObjectLayoutImpl> layout; -class RootShaderObjectLayout : public ShaderObjectLayoutImpl -{ - typedef ShaderObjectLayoutImpl Super; + /// The offset for this entry point's parameters, relative to the starting offset for the program + BindingOffset offset; + }; + + struct Builder : Super::Builder + { + Builder( + RendererBase* renderer, + slang::IComponentType* program, + slang::ProgramLayout* programLayout) + : Super::Builder(renderer, program->getSession()) + , m_program(program) + , m_programLayout(programLayout) + {} + + Result build(RootShaderObjectLayoutImpl** outLayout); + void addGlobalParams(slang::VariableLayoutReflection* globalsLayout); + void addEntryPoint(SlangStage stage, ShaderObjectLayoutImpl* entryPointLayout, slang::EntryPointLayout* slangEntryPoint); + + slang::IComponentType* m_program; + slang::ProgramLayout* m_programLayout; + List<EntryPointInfo> m_entryPoints; + }; + + EntryPointInfo& getEntryPoint(Index index) { return m_entryPoints[index]; } + + List<EntryPointInfo>& getEntryPoints() { return m_entryPoints; } -public: - ~RootShaderObjectLayout(); static Result create( - DeviceImpl* renderer, + RendererBase* renderer, slang::IComponentType* program, slang::ProgramLayout* programLayout, - RootShaderObjectLayout** outLayout); + RootShaderObjectLayoutImpl** outLayout); + + slang::IComponentType* getSlangProgram() const { return m_program; } + slang::ProgramLayout* getSlangProgramLayout() const { return m_programLayout; } + protected: -public: + Result _init(Builder const* builder); + + ComPtr<slang::IComponentType> m_program; + slang::ProgramLayout* m_programLayout = nullptr; + + List<EntryPointInfo> m_entryPoints; }; } // namespace metal diff --git a/tools/gfx/metal/metal-shader-object.cpp b/tools/gfx/metal/metal-shader-object.cpp index d50c4021a..fde04f9ba 100644 --- a/tools/gfx/metal/metal-shader-object.cpp +++ b/tools/gfx/metal/metal-shader-object.cpp @@ -1,7 +1,8 @@ // metal-shader-object.cpp +#include "metal-shader-object.h" +#include "metal-sampler.h" -#include "metal-command-buffer.h" -#include "metal-command-encoder.h" +#include "metal-device.h" namespace gfx { @@ -12,7 +13,9 @@ namespace metal { Result ShaderObjectImpl::create( - IDevice* device, ShaderObjectLayoutImpl* layout, ShaderObjectImpl** outShaderObject) + IDevice* device, + ShaderObjectLayoutImpl* layout, + ShaderObjectImpl** outShaderObject) { auto object = RefPtr<ShaderObjectImpl>(new ShaderObjectImpl()); SLANG_RETURN_ON_FAIL(object->init(device, layout)); @@ -21,88 +24,559 @@ Result ShaderObjectImpl::create( return SLANG_OK; } -RendererBase* ShaderObjectImpl::getDevice() { return m_layout->getDevice(); } - -GfxCount ShaderObjectImpl::getEntryPointCount() { return 0; } - -Result ShaderObjectImpl::getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) +ShaderObjectImpl::~ShaderObjectImpl() { - *outEntryPoint = nullptr; - return SLANG_E_NOT_IMPLEMENTED; } -const void* ShaderObjectImpl::getRawData() { return m_data.getBuffer(); } +SLANG_NO_THROW Result SLANG_MCALL + ShaderObjectImpl::setData(ShaderOffset const& inOffset, void const* data, size_t inSize) +{ + Index offset = inOffset.uniformOffset; + Index size = inSize; -Size ShaderObjectImpl::getSize() { return (Size)m_data.getCount(); } + char* dest = m_data.getBuffer(); + Index availableSize = m_data.getCount(); -// TODO: Change size_t and Index to Size? -Result ShaderObjectImpl::setData(ShaderOffset const& inOffset, void const* data, size_t inSize) -{ - return SLANG_E_NOT_IMPLEMENTED; -} + // TODO: We really should bounds-check access rather than silently ignoring sets + // that are too large, but we have several test cases that set more data than + // an object actually stores on several targets... + // + if (offset < 0) + { + size += offset; + offset = 0; + } + if ((offset + size) >= availableSize) + { + size = availableSize - offset; + } -Result ShaderObjectImpl::setResource(ShaderOffset const& offset, IResourceView* resourceView) -{ - return SLANG_E_NOT_IMPLEMENTED; + memcpy(dest + offset, data, size); + + m_isConstantBufferDirty = true; + + return SLANG_OK; } -Result ShaderObjectImpl::setSampler(ShaderOffset const& offset, ISamplerState* sampler) +SLANG_NO_THROW Result SLANG_MCALL + ShaderObjectImpl::setResource(ShaderOffset const& offset, IResourceView* resourceView) { - return SLANG_E_NOT_IMPLEMENTED; + if (offset.bindingRangeIndex < 0) + return SLANG_E_INVALID_ARG; + auto layout = getLayout(); + if (offset.bindingRangeIndex >= layout->getBindingRangeCount()) + return SLANG_E_INVALID_ARG; + auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex); + + auto resourceViewImpl = static_cast<ResourceViewImpl*>(resourceView); + switch (bindingRange.bindingType) + { + case slang::BindingType::Texture: + case slang::BindingType::MutableTexture: + SLANG_ASSERT(resourceViewImpl->m_type == ResourceViewImpl::ViewType::Texture); + m_textures[bindingRange.baseIndex + offset.bindingArrayIndex] = static_cast<TextureResourceViewImpl*>(resourceView); + break; + case slang::BindingType::RawBuffer: + case slang::BindingType::ConstantBuffer: + case slang::BindingType::MutableRawBuffer: + SLANG_ASSERT(resourceViewImpl->m_type == ResourceViewImpl::ViewType::Buffer); + m_buffers[bindingRange.baseIndex + offset.bindingArrayIndex] = static_cast<BufferResourceViewImpl*>(resourceView); + break; + case slang::BindingType::TypedBuffer: + case slang::BindingType::MutableTypedBuffer: + SLANG_ASSERT(!"Not implemented"); + // SLANG_ASSERT(resourceViewImpl->m_type == ResourceViewImpl::ViewType::TexelBuffer); + // m_textures[bindingRange.baseIndex + offset.bindingArrayIndex] = static_cast<TextureResourceViewImpl*>(resourceView); + break; + } + return SLANG_OK; } -Result ShaderObjectImpl::setCombinedTextureSampler( - ShaderOffset const& offset, IResourceView* textureView, ISamplerState* sampler) +SLANG_NO_THROW Result SLANG_MCALL ShaderObjectImpl::setSampler(ShaderOffset const& offset, ISamplerState* sampler) { - return SLANG_E_NOT_IMPLEMENTED; + if (offset.bindingRangeIndex < 0) + return SLANG_E_INVALID_ARG; + auto layout = getLayout(); + if (offset.bindingRangeIndex >= layout->getBindingRangeCount()) + return SLANG_E_INVALID_ARG; + auto& bindingRange = layout->getBindingRange(offset.bindingRangeIndex); + + m_samplers[bindingRange.baseIndex + offset.bindingArrayIndex] = static_cast<SamplerStateImpl*>(sampler); + return SLANG_OK; } Result ShaderObjectImpl::init(IDevice* device, ShaderObjectLayoutImpl* layout) { - return SLANG_E_NOT_IMPLEMENTED; -} + m_layout = layout; + + // If the layout tells us that there is any uniform data, + // then we will allocate a CPU memory buffer to hold that data + // while it is being set from the host. + // + // Once the user is done setting the parameters/fields of this + // shader object, we will produce a GPU-memory version of the + // uniform data (which includes values from this object and + // any existential-type sub-objects). + // + size_t uniformSize = layout->getElementTypeLayout()->getSize(); + if (uniformSize) + { + m_data.setCount(uniformSize); + memset(m_data.getBuffer(), 0, uniformSize); + } + m_buffers.setCount(layout->getBufferCount()); + m_textures.setCount(layout->getTextureCount()); + m_samplers.setCount(layout->getSamplerCount()); + // If the layout specifies that we have any sub-objects, then + // we need to size the array to account for them. + // + Index subObjectCount = layout->getSubObjectCount(); + m_objects.setCount(subObjectCount); -Result EntryPointShaderObject::create( - IDevice* device, EntryPointLayout* layout, EntryPointShaderObject** outShaderObject) + for (auto subObjectRangeInfo : layout->getSubObjectRanges()) + { + auto subObjectLayout = subObjectRangeInfo.layout; + + // In the case where the sub-object range represents an + // existential-type leaf field (e.g., an `IBar`), we + // cannot pre-allocate the object(s) to go into that + // range, since we can't possibly know what to allocate + // at this point. + // + if (!subObjectLayout) + continue; + // + // Otherwise, we will allocate a sub-object to fill + // in each entry in this range, based on the layout + // information we already have. + + auto& bindingRangeInfo = layout->getBindingRange(subObjectRangeInfo.bindingRangeIndex); + for (Index i = 0; i < bindingRangeInfo.count; ++i) + { + RefPtr<ShaderObjectImpl> subObject; + SLANG_RETURN_ON_FAIL( + ShaderObjectImpl::create(device, subObjectLayout, subObject.writeRef())); + m_objects[bindingRangeInfo.subObjectIndex + i] = subObject; + } + } + + return SLANG_OK; +} + +Result ShaderObjectImpl::_writeOrdinaryData( + void* dest, + size_t destSize, + ShaderObjectLayoutImpl* layout) { - RefPtr<EntryPointShaderObject> object = new EntryPointShaderObject(); - SLANG_RETURN_ON_FAIL(object->init(device, layout)); + // We start by simply writing in the ordinary data contained directly in this object. + // + auto src = m_data.getBuffer(); + auto srcSize = size_t(m_data.getCount()); + SLANG_ASSERT(srcSize <= destSize); + memcpy(dest, src, srcSize); + + // In the case where this object has any sub-objects of + // existential/interface type, we need to recurse on those objects + // that need to write their state into an appropriate "pending" allocation. + // + // Note: Any values that could fit into the "payload" included + // in the existential-type field itself will have already been + // written as part of `setObject()`. This loop only needs to handle + // those sub-objects that do not "fit." + // + // An implementers looking at this code might wonder if things could be changed + // so that *all* writes related to sub-objects for interface-type fields could + // be handled in this one location, rather than having some in `setObject()` and + // others handled here. + // + Index subObjectRangeCounter = 0; + for (auto const& subObjectRangeInfo : layout->getSubObjectRanges()) + { + Index subObjectRangeIndex = subObjectRangeCounter++; + auto const& bindingRangeInfo = layout->getBindingRange(subObjectRangeInfo.bindingRangeIndex); + + // We only need to handle sub-object ranges for interface/existential-type fields, + // because fields of constant-buffer or parameter-block type are responsible for + // the ordinary/uniform data of their own existential/interface-type sub-objects. + // + if (bindingRangeInfo.bindingType != slang::BindingType::ExistentialValue) + continue; + + // Each sub-object range represents a single "leaf" field, but might be nested + // under zero or more outer arrays, such that the number of existential values + // in the same range can be one or more. + // + auto count = bindingRangeInfo.count; + + // We are not concerned with the case where the existential value(s) in the range + // git into the payload part of the leaf field. + // + // In the case where the value didn't fit, the Slang layout strategy would have + // considered the requirements of the value as a "pending" allocation, and would + // allocate storage for the ordinary/uniform part of that pending allocation inside + // of the parent object's type layout. + // + // Here we assume that the Slang reflection API can provide us with a single byte + // offset and stride for the location of the pending data allocation in the specialized + // type layout, which will store the values for this sub-object range. + // + // TODO: The reflection API functions we are assuming here haven't been implemented + // yet, so the functions being called here are stubs. + // + // TODO: It might not be that a single sub-object range can reliably map to a single + // contiguous array with a single stride; we need to carefully consider what the layout + // logic does for complex cases with multiple layers of nested arrays and structures. + // + size_t subObjectRangePendingDataOffset = subObjectRangeInfo.offset.pendingOrdinaryData; + size_t subObjectRangePendingDataStride = subObjectRangeInfo.stride.pendingOrdinaryData; + + // If the range doesn't actually need/use the "pending" allocation at all, then + // we need to detect that case and skip such ranges. + // + // TODO: This should probably be handled on a per-object basis by caching a "does it fit?" + // bit as part of the information for bound sub-objects, given that we already + // compute the "does it fit?" status as part of `setObject()`. + // + if (subObjectRangePendingDataOffset == 0) + continue; + + for (Slang::Index i = 0; i < count; ++i) + { + auto subObject = m_objects[bindingRangeInfo.subObjectIndex + i]; + + ShaderObjectLayoutImpl* subObjectLayout = subObject->getLayout(); + + auto subObjectOffset = subObjectRangePendingDataOffset + i * subObjectRangePendingDataStride; + + auto subObjectDest = (char*)dest + subObjectOffset; + + subObject->_writeOrdinaryData(subObjectDest, destSize - subObjectOffset, subObjectLayout); + } + } - returnRefPtrMove(outShaderObject, object); return SLANG_OK; } +Result ShaderObjectImpl::_ensureOrdinaryDataBufferCreatedIfNeeded( + DeviceImpl* device, + ShaderObjectLayoutImpl* layout) +{ + auto ordinaryDataSize = layout->getTotalOrdinaryDataSize(); + if (ordinaryDataSize == 0) + return SLANG_OK; + + // If we have already created a buffer to hold ordinary data, then we should + // simply re-use that buffer rather than re-create it. + if (!m_ordinaryDataBuffer) + { + ComPtr<IBufferResource> bufferResourcePtr; + IBufferResource::Desc bufferDesc = {}; + bufferDesc.type = IResource::Type::Buffer; + bufferDesc.sizeInBytes = ordinaryDataSize; + bufferDesc.defaultState = ResourceState::ConstantBuffer; + bufferDesc.allowedStates = + ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); + bufferDesc.memoryType = MemoryType::Upload; + SLANG_RETURN_ON_FAIL( + device->createBufferResource(bufferDesc, nullptr, bufferResourcePtr.writeRef())); + m_ordinaryDataBuffer = static_cast<BufferResourceImpl*>(bufferResourcePtr.get()); + } + + if (m_isConstantBufferDirty) + { + // Once the buffer is allocated, we can use `_writeOrdinaryData` to fill it in. + // + // Note that `_writeOrdinaryData` is potentially recursive in the case + // where this object contains interface/existential-type fields, so we + // don't need or want to inline it into this call site. + // + + MemoryRange range = { 0, ordinaryDataSize }; + void* ordinaryData; + SLANG_RETURN_ON_FAIL(m_ordinaryDataBuffer->map(&range, &ordinaryData)); + auto result = _writeOrdinaryData(ordinaryData, ordinaryDataSize, layout); + m_ordinaryDataBuffer->unmap(&range); + m_isConstantBufferDirty = false; + return result; + } + return SLANG_OK; +} -Result EntryPointShaderObject::init(IDevice* device, EntryPointLayout* layout) +Result ShaderObjectImpl::_bindOrdinaryDataBufferIfNeeded( + BindingContext* context, + BindingOffset& ioOffset, + ShaderObjectLayoutImpl* layout) { - //SLANG_RETURN_ON_FAIL(Super::init(device, layout)); - return SLANG_E_NOT_IMPLEMENTED; + // We start by ensuring that the buffer is created, if it is needed. + // + SLANG_RETURN_ON_FAIL(_ensureOrdinaryDataBufferCreatedIfNeeded(context->device, layout)); + + // If we did indeed need/create a buffer, then we must bind it + // into root binding state. + // + if (m_ordinaryDataBuffer) + { + context->setBuffer(m_ordinaryDataBuffer->m_buffer.get(), ioOffset.buffer); + ioOffset.buffer++; + } + + return SLANG_OK; } +Result ShaderObjectImpl::bindAsConstantBuffer( + BindingContext* context, + BindingOffset const& inOffset, + ShaderObjectLayoutImpl* layout) +{ + // When binding a `ConstantBuffer<X>` we need to first bind a constant + // buffer for any "ordinary" data in `X`, and then bind the remaining + // resources and sub-objects. + // + // The one important detail to keep track of its that *if* we bind + // a constant buffer for ordinary data we will need to account for + // it in the offset we use for binding the remaining data. That + // detail is dealt with here by the way that `_bindOrdinaryDataBufferIfNeeded` + // will modify the `offset` parameter if it binds anything. + // + BindingOffset offset = inOffset; + SLANG_RETURN_ON_FAIL(_bindOrdinaryDataBufferIfNeeded(context, /*inout*/ offset, layout)); + + // Once the ordinary data buffer is bound, we can move on to binding + // the rest of the state, which can use logic shared with the case + // for interface-type sub-object ranges. + // + // Note that this call will use the `offset` value that might have + // been modified during `_bindOrindaryDataBufferIfNeeded`. + // + SLANG_RETURN_ON_FAIL(bindAsValue(context, offset, layout)); -GfxCount RootShaderObjectImpl::getEntryPointCount() { return (GfxCount)m_entryPoints.getCount(); } + return SLANG_OK; +} -Result RootShaderObjectImpl::getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) +Result ShaderObjectImpl::bindAsValue( + BindingContext* context, + BindingOffset const& offset, + ShaderObjectLayoutImpl* layout) { - returnComPtr(outEntryPoint, m_entryPoints[index]); + // We start by iterating over the binding ranges in this type, isolating + // just those ranges that represent buffers, textures, and samplers. + // In each loop we will bind the values stored for those binding ranges + // to the correct metal resource indices (based on the `registerOffset` field + // stored in the bindinge range). + + for (auto bindingRangeIndex : layout->getBufferRanges()) + { + auto const& bindingRange = layout->getBindingRange(bindingRangeIndex); + auto count = (uint32_t)bindingRange.count; + auto baseIndex = (uint32_t)bindingRange.baseIndex; + auto registerOffset = bindingRange.registerOffset + offset.buffer; + for (uint32_t i = 0; i < count; ++i) + { + auto buffer = m_buffers[baseIndex + i]; + context->setBuffer(buffer ? buffer->m_buffer->m_buffer.get() : nullptr, registerOffset + i); + } + } + + for (auto bindingRangeIndex : layout->getTextureRanges()) + { + auto const& bindingRange = layout->getBindingRange(bindingRangeIndex); + auto count = (uint32_t)bindingRange.count; + auto baseIndex = (uint32_t)bindingRange.baseIndex; + auto registerOffset = bindingRange.registerOffset + offset.texture; + for (uint32_t i = 0; i < count; ++i) + { + auto texture = m_textures[baseIndex + i]; + context->setTexture(texture ? texture->m_texture->m_texture.get() : nullptr, registerOffset + i); + } + } + + for (auto bindingRangeIndex : layout->getSamplerRanges()) + { + auto const& bindingRange = layout->getBindingRange(bindingRangeIndex); + auto count = (uint32_t)bindingRange.count; + auto baseIndex = (uint32_t)bindingRange.baseIndex; + auto registerOffset = bindingRange.registerOffset + offset.sampler; + for (uint32_t i = 0; i < count; ++i) + { + auto sampler = m_samplers[baseIndex + i]; + context->setSampler(sampler ? sampler->m_samplerState.get() : nullptr, registerOffset + i); + } + } + + // Once all the simple binding ranges are dealt with, we will bind + // all of the sub-objects in sub-object ranges. + // + for (auto const& subObjectRange : layout->getSubObjectRanges()) + { + auto subObjectLayout = subObjectRange.layout; + auto const& bindingRange = layout->getBindingRange(subObjectRange.bindingRangeIndex); + Index count = bindingRange.count; + Index subObjectIndex = bindingRange.subObjectIndex; + + // The starting offset for a sub-object range was computed + // from Slang reflection information, so we can apply it here. + // + BindingOffset rangeOffset = offset; + rangeOffset += subObjectRange.offset; + + // Similarly, the "stride" between consecutive objects in + // the range was also pre-computed. + // + BindingOffset rangeStride = subObjectRange.stride; + + switch (bindingRange.bindingType) + { + case slang::BindingType::ConstantBuffer: + case slang::BindingType::ParameterBlock: + { + BindingOffset objOffset = rangeOffset; + for (Index i = 0; i < count; ++i) + { + auto subObject = m_objects[subObjectIndex + i]; + + // Unsurprisingly, we bind each object in the range as + // a constant buffer. + // + subObject->bindAsConstantBuffer(context, objOffset, subObjectLayout); + + objOffset += rangeStride; + } + } + break; + +#if 0 + case slang::BindingType::ExistentialValue: + // We can only bind information for existential-typed sub-object + // ranges if we have a static type that we are able to specialize to. + // + if (subObjectLayout) + { + // The data for objects in this range will always be bound into + // the "pending" allocation for the parent block/buffer/object. + // As a result, the offset for the first object in the range + // will come from the `pending` part of the range's offset. + // + SimpleBindingOffset objOffset = rangeOffset.pending; + SimpleBindingOffset objStride = rangeStride.pending; + + for (Index i = 0; i < count; ++i) + { + auto subObject = m_objects[subObjectIndex + i]; + subObject->bindAsValue(context, BindingOffset(objOffset), subObjectLayout); + + objOffset += objStride; + } + } + break; +#endif + + default: + break; + } + } + return SLANG_OK; } -Result RootShaderObjectImpl::copyFrom(IShaderObject* object, ITransientResourceHeap* transientHeap) +Result RootShaderObjectImpl::create( + IDevice* device, + RootShaderObjectLayoutImpl* layout, + RootShaderObjectImpl** outShaderObject) { - return SLANG_E_NOT_IMPLEMENTED; + RefPtr<RootShaderObjectImpl> object = new RootShaderObjectImpl(); + SLANG_RETURN_ON_FAIL(object->init(device, layout)); + + returnRefPtrMove(outShaderObject, object); + return SLANG_OK; } Result RootShaderObjectImpl::collectSpecializationArgs(ExtendedShaderObjectTypeList& args) { - return SLANG_E_NOT_IMPLEMENTED; + SLANG_RETURN_ON_FAIL(ShaderObjectImpl::collectSpecializationArgs(args)); + for (auto& entryPoint : m_entryPoints) + { + SLANG_RETURN_ON_FAIL(entryPoint->collectSpecializationArgs(args)); + } + return SLANG_OK; } -Result RootShaderObjectImpl::init(IDevice* device, RootShaderObjectLayout* layout) +Result RootShaderObjectImpl::bindAsRoot( + BindingContext* context, + RootShaderObjectLayoutImpl* layout) { - return SLANG_E_NOT_IMPLEMENTED; + // When binding an entire root shader object, we need to deal with + // the way that specialization might have allocated space for "pending" + // parameter data after all the primary parameters. + // + // We start by initializing an offset that will store zeros for the + // primary data, an the computed offset from the specialized layout + // for pending data. + // + BindingOffset offset; +#if 0 + offset.pending = layout->getPendingDataOffset(); +#endif + + // Note: We could *almost* call `bindAsConstantBuffer()` here to bind + // the state of the root object itself, but there is an important + // detail that means we can't: + // + // The `_bindOrdinaryDataBufferIfNeeded` operation automatically + // increments the offset parameter if it binds a buffer, so that + // subsequently bindings will be adjusted. However, the reflection + // information computed for root shader parameters is absolute rather + // than relative to the default constant buffer (if any). + // + // TODO: Quite technically, the ordinary data buffer for the global + // scope is *not* guaranteed to be at offset zero, so this logic should + // really be querying an appropriate absolute offset from `layout`. + // +#if 0 + BindingOffset ordinaryDataBufferOffset = offset; + SLANG_RETURN_ON_FAIL(_bindOrdinaryDataBufferIfNeeded(context, /*inout*/ ordinaryDataBufferOffset, layout)); +#endif + SLANG_RETURN_ON_FAIL(bindAsValue(context, offset, layout)); + + // Once the state stored in the root shader object itself has been bound, + // we turn our attention to the entry points and their parameters. + // + auto entryPointCount = m_entryPoints.getCount(); + for (Index i = 0; i < entryPointCount; ++i) + { + auto entryPoint = m_entryPoints[i]; + auto const& entryPointInfo = layout->getEntryPoint(i); + + // Each entry point will be bound at some offset relative to where + // the root shader parameters start. + // + BindingOffset entryPointOffset = offset; + entryPointOffset += entryPointInfo.offset; + + // An entry point can simply be bound as a constant buffer, because + // the absolute offsets as are used for the global scope do not apply + // (because entry points don't need to deal with explicit bindings). + // + SLANG_RETURN_ON_FAIL(entryPoint->bindAsConstantBuffer(context, entryPointOffset, entryPointInfo.layout)); + } + + return SLANG_OK; +} + +Result RootShaderObjectImpl::init(IDevice* device, RootShaderObjectLayoutImpl* layout) +{ + SLANG_RETURN_ON_FAIL(Super::init(device, layout)); + + for (auto entryPointInfo : layout->getEntryPoints()) + { + RefPtr<ShaderObjectImpl> entryPoint; + SLANG_RETURN_ON_FAIL( + ShaderObjectImpl::create(device, entryPointInfo.layout, entryPoint.writeRef())); + m_entryPoints.add(entryPoint); + } + + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-shader-object.h b/tools/gfx/metal/metal-shader-object.h index cfdf5be41..b9695febe 100644 --- a/tools/gfx/metal/metal-shader-object.h +++ b/tools/gfx/metal/metal-shader-object.h @@ -1,121 +1,182 @@ // metal-shader-object.h #pragma once - +#include "metal-base.h" #include "metal-resource-views.h" #include "metal-sampler.h" #include "metal-shader-object-layout.h" +#include "metal-helper-functions.h" + namespace gfx { -namespace metal -{ +using namespace Slang; -struct CombinedTextureSamplerSlot +namespace metal { - RefPtr<TextureResourceViewImpl> textureView; - RefPtr<SamplerStateImpl> sampler; - operator bool() { return textureView && sampler; } -}; class ShaderObjectImpl - : public ShaderObjectBaseImpl<ShaderObjectImpl, ShaderObjectLayoutImpl, SimpleShaderObjectData> + : public ShaderObjectBaseImpl< + ShaderObjectImpl, + ShaderObjectLayoutImpl, + SimpleShaderObjectData> { public: static Result create( - IDevice* device, ShaderObjectLayoutImpl* layout, ShaderObjectImpl** outShaderObject); + IDevice* device, + ShaderObjectLayoutImpl* layout, + ShaderObjectImpl** outShaderObject); - RendererBase* getDevice(); + ~ShaderObjectImpl(); - virtual SLANG_NO_THROW GfxCount SLANG_MCALL getEntryPointCount() override; + RendererBase* getDevice() { return m_layout->getDevice(); } - virtual SLANG_NO_THROW Result SLANG_MCALL - getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) override; + SLANG_NO_THROW GfxCount SLANG_MCALL getEntryPointCount() SLANG_OVERRIDE { return 0; } - virtual SLANG_NO_THROW const void* SLANG_MCALL getRawData() override; + SLANG_NO_THROW Result SLANG_MCALL getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) + SLANG_OVERRIDE + { + *outEntryPoint = nullptr; + return SLANG_OK; + } - virtual SLANG_NO_THROW Size SLANG_MCALL getSize() override; + virtual SLANG_NO_THROW const void* SLANG_MCALL getRawData() override + { + return m_data.getBuffer(); + } - // TODO: Changed size_t to Size? inSize assigned to an Index variable inside implementation - virtual SLANG_NO_THROW Result SLANG_MCALL - setData(ShaderOffset const& inOffset, void const* data, size_t inSize) override; + virtual SLANG_NO_THROW size_t SLANG_MCALL getSize() override + { + return (size_t)m_data.getCount(); + } - virtual SLANG_NO_THROW Result SLANG_MCALL - setResource(ShaderOffset const& offset, IResourceView* resourceView) override; + SLANG_NO_THROW Result SLANG_MCALL + setData(ShaderOffset const& inOffset, void const* data, size_t inSize) SLANG_OVERRIDE; - virtual SLANG_NO_THROW Result SLANG_MCALL - setSampler(ShaderOffset const& offset, ISamplerState* sampler) override; + SLANG_NO_THROW Result SLANG_MCALL + setResource(ShaderOffset const& offset, IResourceView* resourceView) SLANG_OVERRIDE; - virtual SLANG_NO_THROW Result SLANG_MCALL setCombinedTextureSampler( - ShaderOffset const& offset, IResourceView* textureView, ISamplerState* sampler) override; - -protected: - friend class RootShaderObjectLayout; + SLANG_NO_THROW Result SLANG_MCALL setSampler(ShaderOffset const& offset, ISamplerState* sampler) + SLANG_OVERRIDE; - Result init(IDevice* device, ShaderObjectLayoutImpl* layout); + SLANG_NO_THROW Result SLANG_MCALL setCombinedTextureSampler( + ShaderOffset const& offset, IResourceView* textureView, ISamplerState* sampler) SLANG_OVERRIDE + { + return SLANG_E_NOT_IMPLEMENTED; + } public: -}; -class MutableShaderObjectImpl - : public MutableShaderObject<MutableShaderObjectImpl, ShaderObjectLayoutImpl> -{ -public: -}; -class EntryPointShaderObject : public ShaderObjectImpl -{ - typedef ShaderObjectImpl Super; +protected: + friend class ProgramVars; -public: - static Result create( - IDevice* device, EntryPointLayout* layout, EntryPointShaderObject** outShaderObject); + Result init(IDevice* device, ShaderObjectLayoutImpl* layout); - EntryPointLayout* getLayout(); + /// Write the uniform/ordinary data of this object into the given `dest` buffer at the given `offset` + Result _writeOrdinaryData( + void* dest, + size_t destSize, + ShaderObjectLayoutImpl* layout); + + /// Ensure that the `m_ordinaryDataBuffer` has been created, if it is needed + /// + /// The `layout` type must represent a specialized layout for this + /// type that includes any "pending" data. + /// + Result _ensureOrdinaryDataBufferCreatedIfNeeded( + DeviceImpl* device, + ShaderObjectLayoutImpl* layout); + + /// Bind the buffer for ordinary/uniform data, if needed + /// + /// The `ioOffset` parameter will be updated to reflect the constant buffer + /// register consumed by the ordinary data buffer, if one was bound. + /// + Result _bindOrdinaryDataBufferIfNeeded( + BindingContext* context, + BindingOffset& ioOffset, + ShaderObjectLayoutImpl* layout); -protected: - Result init(IDevice* device, EntryPointLayout* layout); +public: + /// Bind this object as if it was declared as a `ConstantBuffer<T>` in Slang + Result bindAsConstantBuffer( + BindingContext* context, + BindingOffset const& inOffset, + ShaderObjectLayoutImpl* layout); + + /// Bind this object as a value that appears in the body of another object. + /// + /// This case is directly used when binding an object for an interface-type + /// sub-object range when static specialization is used. It is also used + /// indirectly when binding sub-objects to constant buffer or parameter + /// block ranges. + /// + Result bindAsValue( + BindingContext* context, + BindingOffset const& offset, + ShaderObjectLayoutImpl* layout); + + // Because the binding ranges have already been reflected + // and organized as part of each shader object layout, + // the object itself can store its data in a small number + // of simple arrays. + + /// The buffers that are part of the state of this object + List<RefPtr<BufferResourceViewImpl>> m_buffers; + + /// The textures that are part of the state of this object + List<RefPtr<TextureResourceViewImpl>> m_textures; + + /// The samplers that are part of the state of this object + List<RefPtr<SamplerStateImpl>> m_samplers; + + /// A constant buffer used to stored ordinary data for this object + /// and existential-type sub-objects. + /// + /// Created on demand with `_createOrdinaryDataBufferIfNeeded()` + RefPtr<BufferResourceImpl> m_ordinaryDataBuffer; + + bool m_isConstantBufferDirty = true; }; +class MutableShaderObjectImpl + : public MutableShaderObject< + MutableShaderObjectImpl, + ShaderObjectLayoutImpl> +{}; class RootShaderObjectImpl : public ShaderObjectImpl { - using Super = ShaderObjectImpl; + typedef ShaderObjectImpl Super; public: - // Override default reference counting behavior to disable lifetime management. - // Root objects are managed by command buffer and does not need to be freed by the user. virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } -public: - RootShaderObjectLayout* getLayout(); + static Result create(IDevice* device, RootShaderObjectLayoutImpl* layout, RootShaderObjectImpl** outShaderObject); - RootShaderObjectLayout* getSpecializedLayout(); + Result init(IDevice* device, RootShaderObjectLayoutImpl* layout); - List<RefPtr<EntryPointShaderObject>> const& getEntryPoints() const; + RootShaderObjectLayoutImpl* getLayout() { return static_cast<RootShaderObjectLayoutImpl*>(m_layout.Ptr()); } - virtual GfxCount SLANG_MCALL getEntryPointCount() override; - virtual Result SLANG_MCALL getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) override; + GfxCount SLANG_MCALL getEntryPointCount() SLANG_OVERRIDE { return (GfxCount)m_entryPoints.getCount(); } + SlangResult SLANG_MCALL getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) SLANG_OVERRIDE + { + returnComPtr(outEntryPoint, m_entryPoints[index]); + return SLANG_OK; + } - virtual SLANG_NO_THROW Result SLANG_MCALL - copyFrom(IShaderObject* object, ITransientResourceHeap* transientHeap) override; + virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override; -#if 0 /// Bind this object as a root shader object Result bindAsRoot( - PipelineCommandEncoder* encoder, - RootBindingContext& context, - RootShaderObjectLayout* layout); -#endif - - virtual Result collectSpecializationArgs(ExtendedShaderObjectTypeList& args) override; + BindingContext* context, + RootShaderObjectLayoutImpl* specializedLayout); -public: - Result init(IDevice* device, RootShaderObjectLayout* layout); - List<RefPtr<EntryPointShaderObject>> m_entryPoints; +protected: + List<RefPtr<ShaderObjectImpl>> m_entryPoints; }; - } // namespace metal } // namespace gfx diff --git a/tools/gfx/metal/metal-shader-program.cpp b/tools/gfx/metal/metal-shader-program.cpp index 05f5f9b53..1b998bc57 100644 --- a/tools/gfx/metal/metal-shader-program.cpp +++ b/tools/gfx/metal/metal-shader-program.cpp @@ -1,7 +1,7 @@ // metal-shader-program.cpp #include "metal-shader-program.h" - #include "metal-device.h" +#include "metal-util.h" namespace gfx { @@ -20,23 +20,16 @@ ShaderProgramImpl::~ShaderProgramImpl() { } -void ShaderProgramImpl::comFree() { } - -Result ShaderProgramImpl::createShaderModule( - slang::EntryPointReflection* entryPointInfo, ComPtr<ISlangBlob> kernelCode) +Result ShaderProgramImpl::createShaderModule(slang::EntryPointReflection* entryPointInfo, ComPtr<ISlangBlob> kernelCode) { - if (entryPointInfo == nullptr || kernelCode == nullptr || kernelCode->getBufferSize() == 0) - { - return SLANG_E_INVALID_ARG; - } - - auto realEntryPointName = entryPointInfo->getNameOverride(); - std::string sourceStr(static_cast<const char*>(kernelCode->getBufferPointer()), kernelCode->getBufferSize()); - NS::String *nsSourceString = NS::String::alloc()->init(sourceStr.c_str(), NS::UTF8StringEncoding); + m_codeBlobs.add(kernelCode); + const char* realEntryPointName = entryPointInfo->getNameOverride(); + dispatch_data_t data = dispatch_data_create(kernelCode->getBufferPointer(), kernelCode->getBufferSize(), dispatch_get_main_queue(), NULL); NS::Error* error; - MTL::Library* library = m_device->m_device->newLibrary(nsSourceString, nullptr, &error); - if (library == nullptr) + NS::SharedPtr<MTL::Library> library = NS::TransferPtr(m_device->m_device->newLibrary(data, &error)); + if (!library) { + // TODO use better mechanism for reporting errors std::cout << error->localizedDescription()->utf8String() << std::endl; return SLANG_E_INVALID_ARG; } diff --git a/tools/gfx/metal/metal-shader-program.h b/tools/gfx/metal/metal-shader-program.h index 3846328c0..50f568321 100644 --- a/tools/gfx/metal/metal-shader-program.h +++ b/tools/gfx/metal/metal-shader-program.h @@ -15,21 +15,17 @@ namespace metal class ShaderProgramImpl : public ShaderProgramBase { public: - ShaderProgramImpl(DeviceImpl* device); - - ~ShaderProgramImpl(); - - virtual void comFree() override; - BreakableReference<DeviceImpl> m_device; List<String> m_entryPointNames; List<ComPtr<ISlangBlob>> m_codeBlobs; //< To keep storage of code in scope - List<MTL::Library*> m_modules; - RefPtr<RootShaderObjectLayout> m_rootObjectLayout; + List<NS::SharedPtr<MTL::Library>> m_modules; + RefPtr<RootShaderObjectLayoutImpl> m_rootObjectLayout; + + ShaderProgramImpl(DeviceImpl* device); + ~ShaderProgramImpl(); - virtual Result createShaderModule( - slang::EntryPointReflection* entryPointInfo, ComPtr<ISlangBlob> kernelCode) override; + virtual Result createShaderModule(slang::EntryPointReflection* entryPointInfo, ComPtr<ISlangBlob> kernelCode) override; }; diff --git a/tools/gfx/metal/metal-swap-chain.cpp b/tools/gfx/metal/metal-swap-chain.cpp index c80ee8cf8..0a28b214c 100644 --- a/tools/gfx/metal/metal-swap-chain.cpp +++ b/tools/gfx/metal/metal-swap-chain.cpp @@ -19,87 +19,44 @@ ISwapchain* SwapchainImpl::getInterface(const Guid& guid) return nullptr; } -void SwapchainImpl::destroySwapchainAndImages() -{ - m_images.clear(); -} - void SwapchainImpl::getWindowSize(int& widthOut, int& heightOut) const { CocoaUtil::getNSWindowContentSize((void*)m_windowHandle.handleValues[0], &widthOut, &heightOut); } -Result SwapchainImpl::createSwapchainAndImages() -{ - getWindowSize(m_desc.width, m_desc.height); - // Note that we do not actually create/assign textures here, as metal requires that one do so JIT, - // rather than ahead of time. - //m_drawables.setCount(m_desc.imageCount); - for (GfxIndex i = 0; i < m_desc.imageCount; i++) - { - //CA::MetalDrawable* drawable = m_metalLayer->nextDrawable(); - //if (drawable == nullptr) - //{ - // assert(drawable); - // } - // m_drawables[i] = drawable; - //MTL::Texture* tex = drawable->texture(); - - ITextureResource::Desc imageDesc = {}; - imageDesc.allowedStates = ResourceStateSet( - ResourceState::Present, ResourceState::RenderTarget, ResourceState::CopyDestination); - imageDesc.type = IResource::Type::Texture2D; - imageDesc.arraySize = 0; - imageDesc.format = m_desc.format; - imageDesc.size.width = m_desc.width; - imageDesc.size.height = m_desc.height; - imageDesc.size.depth = 1; - imageDesc.numMipLevels = 1; - imageDesc.defaultState = ResourceState::Present; - RefPtr<TextureResourceImpl> image = new TextureResourceImpl(imageDesc, m_renderer); - //image->m_texture = tex; - image->m_texture = nullptr; - m_images.add(image); - } - return SLANG_OK; -} - SwapchainImpl::~SwapchainImpl() { - destroySwapchainAndImages(); - CocoaUtil::destroyMetalLayer(m_renderer->m_metalLayer); + m_images.clear(); + CocoaUtil::destroyMetalLayer(m_metalLayer); } -Result SwapchainImpl::init(DeviceImpl* renderer, const ISwapchain::Desc& desc, WindowHandle window) +Result SwapchainImpl::init(DeviceImpl* device, const ISwapchain::Desc& desc, WindowHandle window) { - m_renderer = renderer; - m_queue = static_cast<CommandQueueImpl*>(desc.queue); - m_windowHandle = window; - m_metalFormat = MetalUtil::getMetalPixelFormat(desc.format); - - int width, height; - getWindowSize(width, height); - CGSize windowSize = {(float)width, (float)height}; - + m_device = device; m_desc = desc; + m_windowHandle = window; + m_metalFormat = MetalUtil::translatePixelFormat(desc.format); + m_currentImageIndex = 0; - m_renderer->m_metalLayer = (CA::MetalLayer*)CocoaUtil::createMetalLayer((void*)window.handleValues[0]); - m_renderer->m_metalLayer->setPixelFormat(m_metalFormat); - m_renderer->m_metalLayer->setDevice(renderer->m_device); - m_renderer->m_metalLayer->setDrawableSize(windowSize); - m_renderer->m_metalLayer->setFramebufferOnly(true); + getWindowSize(m_desc.width, m_desc.height); - createSwapchainAndImages(); + m_metalLayer = (CA::MetalLayer*)CocoaUtil::createMetalLayer((void*)window.handleValues[0]); + if (!m_metalLayer) + { + return SLANG_FAIL; + } + m_metalLayer->setPixelFormat(m_metalFormat); + m_metalLayer->setDevice(m_device->m_device.get()); + m_metalLayer->setDrawableSize(CGSize{(float)m_desc.width, (float)m_desc.height}); + m_metalLayer->setFramebufferOnly(true); return SLANG_OK; } Result SwapchainImpl::getImage(GfxIndex index, ITextureResource** outResource) { - if (m_images.getCount() <= (Index)index) + if (index < 0 || index != m_currentImageIndex) return SLANG_FAIL; - // TODO: iff index == current - m_images[index]->m_isCurrentDrawable = true; returnComPtr(outResource, m_images[index]); return SLANG_OK; } @@ -108,33 +65,83 @@ Result SwapchainImpl::resize(GfxCount width, GfxCount height) { SLANG_UNUSED(width); SLANG_UNUSED(height); - destroySwapchainAndImages(); - return createSwapchainAndImages(); + m_images.clear(); + m_currentImageIndex = -1; + m_currentDrawable.reset(); + getWindowSize(m_desc.width, m_desc.height); + m_metalLayer->setDrawableSize(CGSize{(float)m_desc.width, (float)m_desc.height}); + return SLANG_OK; } Result SwapchainImpl::present() { - // TODO: Expose controls via some other means - static uint32_t frameCount = 0; - static uint32_t maxFrameCount = 32; - ++frameCount; - if (m_renderer->captureEnabled() && frameCount == maxFrameCount) + if (!m_currentDrawable) { - MTL::CaptureManager* captureManager = MTL::CaptureManager::sharedCaptureManager(); - captureManager->stopCapture(); - exit(1); + return SLANG_FAIL; } + + MTL::CommandBuffer* commandBuffer = m_device->m_commandQueue->commandBuffer(); + commandBuffer->presentDrawable(m_currentDrawable.get()); + commandBuffer->commit(); + m_currentDrawable.reset(); return SLANG_OK; + + // // TODO: Expose controls via some other means + // static uint32_t frameCount = 0; + // static uint32_t maxFrameCount = 32; + // ++frameCount; + // if (m_device->captureEnabled() && frameCount == maxFrameCount) + // { + // MTL::CaptureManager* captureManager = MTL::CaptureManager::sharedCaptureManager(); + // captureManager->stopCapture(); + // exit(1); + // } + // return SLANG_OK; } int SwapchainImpl::acquireNextImage() { - // TODO: hardcoded 0 - CA::MetalDrawable* d = m_renderer->m_metalLayer->nextDrawable(); - m_images[0]->m_texture = d->texture(); - m_renderer->m_drawable = d; + AUTORELEASEPOOL + + CA::MetalDrawable* drawable = m_metalLayer->nextDrawable(); + if (drawable == nullptr) + { + return -1; + } + + m_currentDrawable = NS::RetainPtr(drawable); + MTL::Texture* texture = drawable->texture(); + + // Check if we got a texture we've seen before. + for (Index i = 0; i < m_images.getCount(); i++) + { + if (m_images[i]->m_texture.get() == texture) + { + m_currentImageIndex = i; + return m_currentImageIndex; + } + } - return 0; + // Create a new texture object to wrap the drawable's texture. + + ITextureResource::Desc desc = {}; + desc.allowedStates = ResourceStateSet( + ResourceState::Present, ResourceState::RenderTarget, ResourceState::CopyDestination); + desc.type = IResource::Type::Texture2D; + desc.arraySize = 0; + desc.format = m_desc.format; // TODO use actual pixelformat + desc.size.width = texture->width(); + desc.size.height = texture->height(); + desc.size.depth = 1; + desc.numMipLevels = 1; + desc.defaultState = ResourceState::Present; + RefPtr<TextureResourceImpl> image = new TextureResourceImpl(desc, m_device); + image->m_texture = NS::RetainPtr(texture); + + m_currentImageIndex = m_images.getCount(); + m_images.add(image); + + return m_currentImageIndex; } Result SwapchainImpl::setFullScreenMode(bool mode) { return SLANG_FAIL; } diff --git a/tools/gfx/metal/metal-swap-chain.h b/tools/gfx/metal/metal-swap-chain.h index c96a5fd5e..bed1c0f5b 100644 --- a/tools/gfx/metal/metal-swap-chain.h +++ b/tools/gfx/metal/metal-swap-chain.h @@ -23,25 +23,21 @@ public: ISwapchain* getInterface(const Guid& guid); public: + RefPtr<DeviceImpl> m_device; ISwapchain::Desc m_desc; - RefPtr<CommandQueueImpl> m_queue; - ShortList<RefPtr<TextureResourceImpl>> m_images; - ShortList<MTL::Drawable*> m_drawables; - RefPtr<DeviceImpl> m_renderer; - uint32_t m_currentImageIndex = 0; WindowHandle m_windowHandle; + CA::MetalLayer* m_metalLayer = nullptr; + ShortList<RefPtr<TextureResourceImpl>> m_images; + NS::SharedPtr<MTL::Drawable> m_currentDrawable; + Index m_currentImageIndex = -1; MTL::PixelFormat m_metalFormat = MTL::PixelFormat::PixelFormatInvalid; - void destroySwapchainAndImages(); - void getWindowSize(int& widthOut, int& heightOut) const; - Result createSwapchainAndImages(); - public: ~SwapchainImpl(); - Result init(DeviceImpl* renderer, const ISwapchain::Desc& desc, WindowHandle window); + Result init(DeviceImpl* device, const ISwapchain::Desc& desc, WindowHandle window); virtual SLANG_NO_THROW const Desc& SLANG_MCALL getDesc() override { return m_desc; } virtual SLANG_NO_THROW Result SLANG_MCALL diff --git a/tools/gfx/metal/metal-texture.cpp b/tools/gfx/metal/metal-texture.cpp index 092a028b0..81625aa7e 100644 --- a/tools/gfx/metal/metal-texture.cpp +++ b/tools/gfx/metal/metal-texture.cpp @@ -1,5 +1,6 @@ // metal-texture.cpp #include "metal-texture.h" +#include "metal-util.h" namespace gfx { @@ -21,18 +22,20 @@ TextureResourceImpl::~TextureResourceImpl() Result TextureResourceImpl::getNativeResourceHandle(InteropHandle* outHandle) { outHandle->api = InteropHandleAPI::Metal; - outHandle->handleValue = reinterpret_cast<intptr_t>(m_texture); + outHandle->handleValue = reinterpret_cast<intptr_t>(m_texture.get()); return SLANG_OK; } Result TextureResourceImpl::getSharedHandle(InteropHandle* outHandle) { - return SLANG_E_NOT_IMPLEMENTED; + return SLANG_E_NOT_AVAILABLE; } Result TextureResourceImpl::setDebugName(const char* name) { - return SLANG_E_NOT_IMPLEMENTED; + Parent::setDebugName(name); + m_texture->setLabel(MetalUtil::createString(name).get()); + return SLANG_OK; } } // namespace metal diff --git a/tools/gfx/metal/metal-texture.h b/tools/gfx/metal/metal-texture.h index ada0f77aa..853af7ab0 100644 --- a/tools/gfx/metal/metal-texture.h +++ b/tools/gfx/metal/metal-texture.h @@ -16,17 +16,17 @@ class TextureResourceImpl : public TextureResource { public: typedef TextureResource Parent; + TextureResourceImpl(const Desc& desc, DeviceImpl* device); ~TextureResourceImpl(); - MTL::Texture* m_texture = nullptr; - //MTL::PixelFormat m_metalFormat = MTL::PixelFormat::PixelFormatInvalid; - //bool m_isWeakImageReference = false; - bool m_isCurrentDrawable = false; RefPtr<DeviceImpl> m_device; + NS::SharedPtr<MTL::Texture> m_texture; + // TODO still needed? + // MTL::PixelFormat m_metalFormat = MTL::PixelFormat::PixelFormatInvalid; + // bool m_isWeakImageReference = false; - virtual SLANG_NO_THROW Result SLANG_MCALL - getNativeResourceHandle(InteropHandle* outHandle) override; + virtual SLANG_NO_THROW Result SLANG_MCALL getNativeResourceHandle(InteropHandle* outHandle) override; virtual SLANG_NO_THROW Result SLANG_MCALL getSharedHandle(InteropHandle* outHandle) override; diff --git a/tools/gfx/metal/metal-transient-heap.cpp b/tools/gfx/metal/metal-transient-heap.cpp new file mode 100644 index 000000000..8810897dd --- /dev/null +++ b/tools/gfx/metal/metal-transient-heap.cpp @@ -0,0 +1,44 @@ +// metal-transient-heap.cpp +#include "metal-transient-heap.h" + +#include "metal-device.h" +#include "metal-util.h" + +namespace gfx +{ + +using namespace Slang; + +namespace metal +{ + +Result TransientResourceHeapImpl::init(const ITransientResourceHeap::Desc& desc, DeviceImpl* device) +{ + Super::init( + desc, + 256, // TODO + device); + + return SLANG_OK; +} + +TransientResourceHeapImpl::~TransientResourceHeapImpl() +{ +} + +Result TransientResourceHeapImpl::createCommandBuffer(ICommandBuffer** outCmdBuffer) +{ + RefPtr<CommandBufferImpl> commandBuffer = new CommandBufferImpl(); + SLANG_RETURN_ON_FAIL(commandBuffer->init(m_device, this)); + returnComPtr(outCmdBuffer, commandBuffer); + return SLANG_OK; +} + +Result TransientResourceHeapImpl::synchronizeAndReset() +{ + Super::reset(); + return SLANG_OK; +} + +} // namespace metal +} // namespace gfx diff --git a/tools/gfx/metal/metal-transient-heap.h b/tools/gfx/metal/metal-transient-heap.h new file mode 100644 index 000000000..4b3f6dee9 --- /dev/null +++ b/tools/gfx/metal/metal-transient-heap.h @@ -0,0 +1,35 @@ +// metal-transient-heap.h +#pragma once + +#include "metal-base.h" +#include "metal-buffer.h" +#include "metal-command-buffer.h" + +namespace gfx +{ + +using namespace Slang; + +namespace metal +{ + +class TransientResourceHeapImpl + : public TransientResourceHeapBaseImpl<DeviceImpl, BufferResourceImpl> +{ +private: + typedef TransientResourceHeapBaseImpl<DeviceImpl, BufferResourceImpl> Super; + +public: + NS::SharedPtr<MTL::CommandQueue> m_commandQueue; + + Result init(const ITransientResourceHeap::Desc& desc, DeviceImpl* device); + ~TransientResourceHeapImpl(); + +public: + virtual SLANG_NO_THROW Result SLANG_MCALL + createCommandBuffer(ICommandBuffer** outCommandBuffer) override; + virtual SLANG_NO_THROW Result SLANG_MCALL synchronizeAndReset() override; +}; + +} // namespace metal +} // namespace gfx diff --git a/tools/gfx/metal/metal-util.cpp b/tools/gfx/metal/metal-util.cpp index a8c8da6b6..d7d2a793d 100644 --- a/tools/gfx/metal/metal-util.cpp +++ b/tools/gfx/metal/metal-util.cpp @@ -7,174 +7,255 @@ namespace gfx { -using namespace MTL; +MTL::PixelFormat MetalUtil::translatePixelFormat(Format format) +{ + switch (format) + { + case Format::R32G32B32A32_TYPELESS: return MTL::PixelFormatRGBA32Float; + case Format::R32G32B32_TYPELESS: return MTL::PixelFormatInvalid; + case Format::R32G32_TYPELESS: return MTL::PixelFormatRG32Float; + case Format::R32_TYPELESS: return MTL::PixelFormatR32Float; + + case Format::R16G16B16A16_TYPELESS: return MTL::PixelFormatRGBA16Float; + case Format::R16G16_TYPELESS: return MTL::PixelFormatRG16Float; + case Format::R16_TYPELESS: return MTL::PixelFormatR16Float; + + case Format::R8G8B8A8_TYPELESS: return MTL::PixelFormatRGBA8Unorm; + case Format::R8G8_TYPELESS: return MTL::PixelFormatRG8Unorm; + case Format::R8_TYPELESS: return MTL::PixelFormatR8Unorm; + case Format::B8G8R8A8_TYPELESS: return MTL::PixelFormatBGRA8Unorm; + + case Format::R32G32B32A32_FLOAT: return MTL::PixelFormatRGBA32Float; + case Format::R32G32B32_FLOAT: return MTL::PixelFormatInvalid; + case Format::R32G32_FLOAT: return MTL::PixelFormatRG32Float; + case Format::R32_FLOAT: return MTL::PixelFormatR32Float; + + case Format::R16G16B16A16_FLOAT: return MTL::PixelFormatRGBA16Float; + case Format::R16G16_FLOAT: return MTL::PixelFormatRG16Float; + case Format::R16_FLOAT: return MTL::PixelFormatR16Float; + + case Format::R32G32B32A32_UINT: return MTL::PixelFormatRGBA32Uint; + case Format::R32G32B32_UINT: return MTL::PixelFormatInvalid; + case Format::R32G32_UINT: return MTL::PixelFormatRG32Uint; + case Format::R32_UINT: return MTL::PixelFormatR32Uint; + + case Format::R16G16B16A16_UINT: return MTL::PixelFormatRGBA16Uint; + case Format::R16G16_UINT: return MTL::PixelFormatRG16Uint; + case Format::R16_UINT: return MTL::PixelFormatR16Uint; + + case Format::R8G8B8A8_UINT: return MTL::PixelFormatRGBA8Uint; + case Format::R8G8_UINT: return MTL::PixelFormatRG8Uint; + case Format::R8_UINT: return MTL::PixelFormatR8Uint; + + case Format::R32G32B32A32_SINT: return MTL::PixelFormatRGBA32Sint; + case Format::R32G32B32_SINT: return MTL::PixelFormatInvalid; + case Format::R32G32_SINT: return MTL::PixelFormatRG32Sint; + case Format::R32_SINT: return MTL::PixelFormatR32Sint; + + case Format::R16G16B16A16_SINT: return MTL::PixelFormatRGBA16Sint; + case Format::R16G16_SINT: return MTL::PixelFormatRG16Sint; + case Format::R16_SINT: return MTL::PixelFormatR16Sint; + + case Format::R8G8B8A8_SINT: return MTL::PixelFormatRGBA8Sint; + case Format::R8G8_SINT: return MTL::PixelFormatRG8Sint; + case Format::R8_SINT: return MTL::PixelFormatR8Sint; + + case Format::R16G16B16A16_UNORM: return MTL::PixelFormatRGBA16Unorm; + case Format::R16G16_UNORM: return MTL::PixelFormatRG16Unorm; + case Format::R16_UNORM: return MTL::PixelFormatR16Unorm; + + case Format::R8G8B8A8_UNORM: return MTL::PixelFormatRGBA8Unorm; + case Format::R8G8B8A8_UNORM_SRGB: return MTL::PixelFormatRGBA8Unorm_sRGB; + case Format::R8G8_UNORM: return MTL::PixelFormatRG8Unorm; + case Format::R8_UNORM: return MTL::PixelFormatR8Unorm; + case Format::B8G8R8A8_UNORM: return MTL::PixelFormatBGRA8Unorm; + case Format::B8G8R8A8_UNORM_SRGB: return MTL::PixelFormatBGRA8Unorm_sRGB; + case Format::B8G8R8X8_UNORM: return MTL::PixelFormatInvalid; + case Format::B8G8R8X8_UNORM_SRGB: return MTL::PixelFormatInvalid; + + case Format::R16G16B16A16_SNORM: return MTL::PixelFormatRGBA16Snorm; + case Format::R16G16_SNORM: return MTL::PixelFormatRG16Snorm; + case Format::R16_SNORM: return MTL::PixelFormatR16Snorm; + + case Format::R8G8B8A8_SNORM: return MTL::PixelFormatRGBA8Snorm; + case Format::R8G8_SNORM: return MTL::PixelFormatRG8Snorm; + case Format::R8_SNORM: return MTL::PixelFormatR8Snorm; + + case Format::D32_FLOAT: return MTL::PixelFormatDepth32Float; + case Format::D16_UNORM: return MTL::PixelFormatDepth16Unorm; + case Format::D32_FLOAT_S8_UINT: return MTL::PixelFormatDepth32Float_Stencil8; + case Format::R32_FLOAT_X32_TYPELESS: return MTL::PixelFormatInvalid; + + case Format::B4G4R4A4_UNORM: return MTL::PixelFormatABGR4Unorm; + case Format::B5G6R5_UNORM: return MTL::PixelFormatB5G6R5Unorm; + case Format::B5G5R5A1_UNORM: return MTL::PixelFormatA1BGR5Unorm; + + case Format::R9G9B9E5_SHAREDEXP: return MTL::PixelFormatRGB9E5Float; + case Format::R10G10B10A2_TYPELESS: return MTL::PixelFormatInvalid; + case Format::R10G10B10A2_UINT: return MTL::PixelFormatRGB10A2Uint; + case Format::R10G10B10A2_UNORM: return MTL::PixelFormatRGB10A2Unorm; + case Format::R11G11B10_FLOAT: return MTL::PixelFormatRG11B10Float; + + case Format::BC1_UNORM: return MTL::PixelFormatBC1_RGBA; + case Format::BC1_UNORM_SRGB: return MTL::PixelFormatBC1_RGBA_sRGB; + case Format::BC2_UNORM: return MTL::PixelFormatBC2_RGBA; + case Format::BC2_UNORM_SRGB: return MTL::PixelFormatBC2_RGBA_sRGB; + case Format::BC3_UNORM: return MTL::PixelFormatBC3_RGBA; + case Format::BC3_UNORM_SRGB: return MTL::PixelFormatBC3_RGBA_sRGB; + case Format::BC4_UNORM: return MTL::PixelFormatBC4_RUnorm; + case Format::BC4_SNORM: return MTL::PixelFormatBC4_RSnorm; + case Format::BC5_UNORM: return MTL::PixelFormatBC5_RGUnorm; + case Format::BC5_SNORM: return MTL::PixelFormatBC5_RGSnorm; + case Format::BC6H_UF16: return MTL::PixelFormatBC6H_RGBUfloat; + case Format::BC6H_SF16: return MTL::PixelFormatBC6H_RGBFloat; + case Format::BC7_UNORM: return MTL::PixelFormatBC7_RGBAUnorm; + case Format::BC7_UNORM_SRGB: return MTL::PixelFormatBC7_RGBAUnorm_sRGB; + + default: return MTL::PixelFormatInvalid; + } +} -MTL::VertexFormat MetalUtil::getMetalVertexFormat(Format format) +MTL::VertexFormat MetalUtil::translateVertexFormat(Format format) { switch (format) { - case Format::R8G8_UINT: return VertexFormatUChar2; + case Format::R8G8_UINT: return MTL::VertexFormatUChar2; // VertexFormatUChar3 - case Format::R8G8B8A8_UINT: return VertexFormatUChar4; - case Format::R8G8_SINT: return VertexFormatChar2; + case Format::R8G8B8A8_UINT: return MTL::VertexFormatUChar4; + case Format::R8G8_SINT: return MTL::VertexFormatChar2; // return VertexFormatChar3 - case Format::R8G8B8A8_SINT: return VertexFormatChar4; - case Format::R8G8_UNORM: return VertexFormatUChar2Normalized; + case Format::R8G8B8A8_SINT: return MTL::VertexFormatChar4; + case Format::R8G8_UNORM: return MTL::VertexFormatUChar2Normalized; // return VertexFormatUChar3Normalized; - case Format::R8G8B8A8_UNORM: return VertexFormatUChar4Normalized; - case Format::R8G8_SNORM: return VertexFormatChar2Normalized; + case Format::R8G8B8A8_UNORM: return MTL::VertexFormatUChar4Normalized; + case Format::R8G8_SNORM: return MTL::VertexFormatChar2Normalized; // return VertexFormatChar3Normalized - case Format::R8G8B8A8_SNORM: return VertexFormatChar4Normalized; - case Format::R16G16_UINT: return VertexFormatUShort2; + case Format::R8G8B8A8_SNORM: return MTL::VertexFormatChar4Normalized; + case Format::R16G16_UINT: return MTL::VertexFormatUShort2; // return VertexFormatUShort3; - case Format::R16G16B16A16_UINT: return VertexFormatUShort4; - case Format::R16G16_SINT: return VertexFormatShort2; + case Format::R16G16B16A16_UINT: return MTL::VertexFormatUShort4; + case Format::R16G16_SINT: return MTL::VertexFormatShort2; // return VertexFormatShort3; - case Format::R16G16B16A16_SINT: return VertexFormatShort4; - case Format::R16G16_UNORM: return VertexFormatUShort2Normalized; + case Format::R16G16B16A16_SINT: return MTL::VertexFormatShort4; + case Format::R16G16_UNORM: return MTL::VertexFormatUShort2Normalized; // return VertexFormatUShort3Normalized; - case Format::R16G16B16A16_UNORM: return VertexFormatUShort4Normalized; - case Format::R16G16_SNORM: return VertexFormatShort2Normalized; + case Format::R16G16B16A16_UNORM: return MTL::VertexFormatUShort4Normalized; + case Format::R16G16_SNORM: return MTL::VertexFormatShort2Normalized; // return VertexFormatShort3Normalized; - case Format::R16G16B16A16_SNORM: return VertexFormatShort4Normalized; - case Format::R16G16_FLOAT: return VertexFormatHalf2; + case Format::R16G16B16A16_SNORM: return MTL::VertexFormatShort4Normalized; + case Format::R16G16_FLOAT: return MTL::VertexFormatHalf2; // return VertexFormatHalf3; - case Format::R16G16B16A16_FLOAT: return VertexFormatHalf4; - case Format::R32_FLOAT: return VertexFormatFloat; - case Format::R32G32_FLOAT: return VertexFormatFloat2; - case Format::R32G32B32_FLOAT: return VertexFormatFloat3; - case Format::R32G32B32A32_FLOAT: return VertexFormatFloat4; - case Format::R32_SINT: return VertexFormatInt; - case Format::R32G32_SINT: return VertexFormatInt2; - case Format::R32G32B32_SINT: return VertexFormatInt3; - case Format::R32G32B32A32_SINT: return VertexFormatInt4; - case Format::R32_UINT: return VertexFormatUInt; - case Format::R32G32_UINT: return VertexFormatUInt2; - case Format::R32G32B32_UINT: return VertexFormatUInt3; - case Format::R32G32B32A32_UINT: return VertexFormatUInt4; + case Format::R16G16B16A16_FLOAT: return MTL::VertexFormatHalf4; + case Format::R32_FLOAT: return MTL::VertexFormatFloat; + case Format::R32G32_FLOAT: return MTL::VertexFormatFloat2; + case Format::R32G32B32_FLOAT: return MTL::VertexFormatFloat3; + case Format::R32G32B32A32_FLOAT: return MTL::VertexFormatFloat4; + case Format::R32_SINT: return MTL::VertexFormatInt; + case Format::R32G32_SINT: return MTL::VertexFormatInt2; + case Format::R32G32B32_SINT: return MTL::VertexFormatInt3; + case Format::R32G32B32A32_SINT: return MTL::VertexFormatInt4; + case Format::R32_UINT: return MTL::VertexFormatUInt; + case Format::R32G32_UINT: return MTL::VertexFormatUInt2; + case Format::R32G32B32_UINT: return MTL::VertexFormatUInt3; + case Format::R32G32B32A32_UINT: return MTL::VertexFormatUInt4; // return VertexFormatInt1010102Normalized; - case Format::R10G10B10A2_UNORM: return VertexFormatUInt1010102Normalized; - case Format::B4G4R4A4_UNORM: return VertexFormatUChar4Normalized_BGRA; - case Format::R8_UINT: return VertexFormatUChar; - case Format::R8_SINT: return VertexFormatChar; - case Format::R8_UNORM: return VertexFormatUCharNormalized; - case Format::R8_SNORM: return VertexFormatCharNormalized; - case Format::R16_UINT: return VertexFormatUShort; - case Format::R16_SINT: return VertexFormatShort; - case Format::R16_UNORM: return VertexFormatUShortNormalized; - case Format::R16_SNORM: return VertexFormatShortNormalized; - case Format::R16_FLOAT: return VertexFormatHalf; - case Format::R11G11B10_FLOAT: return VertexFormatFloatRG11B10; - case Format::R9G9B9E5_SHAREDEXP: return VertexFormatFloatRGB9E5; - default: return VertexFormatInvalid; + case Format::R10G10B10A2_UNORM: return MTL::VertexFormatUInt1010102Normalized; + case Format::B4G4R4A4_UNORM: return MTL::VertexFormatUChar4Normalized_BGRA; + case Format::R8_UINT: return MTL::VertexFormatUChar; + case Format::R8_SINT: return MTL::VertexFormatChar; + case Format::R8_UNORM: return MTL::VertexFormatUCharNormalized; + case Format::R8_SNORM: return MTL::VertexFormatCharNormalized; + case Format::R16_UINT: return MTL::VertexFormatUShort; + case Format::R16_SINT: return MTL::VertexFormatShort; + case Format::R16_UNORM: return MTL::VertexFormatUShortNormalized; + case Format::R16_SNORM: return MTL::VertexFormatShortNormalized; + case Format::R16_FLOAT: return MTL::VertexFormatHalf; + case Format::R11G11B10_FLOAT: return MTL::VertexFormatFloatRG11B10; + case Format::R9G9B9E5_SHAREDEXP: return MTL::VertexFormatFloatRGB9E5; + default: return MTL::VertexFormatInvalid; } } -/* static */MTL::PixelFormat MetalUtil::getMetalPixelFormat(Format format) +MTL::SamplerMinMagFilter MetalUtil::translateSamplerMinMagFilter(TextureFilteringMode mode) { - switch (format) + switch (mode) + { + case TextureFilteringMode::Point: + return MTL::SamplerMinMagFilterNearest; + case TextureFilteringMode::Linear: + return MTL::SamplerMinMagFilterLinear; + default: + return MTL::SamplerMinMagFilter(0); + } +} + +MTL::SamplerMipFilter MetalUtil::translateSamplerMipFilter(TextureFilteringMode mode) +{ + switch (mode) + { + case TextureFilteringMode::Point: + return MTL::SamplerMipFilterNearest; + case TextureFilteringMode::Linear: + return MTL::SamplerMipFilterLinear; + default: + return MTL::SamplerMipFilter(0); + } +} + +MTL::SamplerAddressMode MetalUtil::translateSamplerAddressMode(TextureAddressingMode mode) +{ + switch (mode) + { + case TextureAddressingMode::Wrap: + return MTL::SamplerAddressModeRepeat; + case TextureAddressingMode::ClampToEdge: + return MTL::SamplerAddressModeClampToEdge; + case TextureAddressingMode::ClampToBorder: + return MTL::SamplerAddressModeClampToBorderColor; + case TextureAddressingMode::MirrorRepeat: + return MTL::SamplerAddressModeMirrorRepeat; + case TextureAddressingMode::MirrorOnce: + return MTL::SamplerAddressModeMirrorClampToEdge; + default: + return MTL::SamplerAddressMode(0); + } +} + +MTL::CompareFunction MetalUtil::translateCompareFunction(ComparisonFunc func) +{ + switch (func) + { + case ComparisonFunc::Never: + return MTL::CompareFunctionNever; + case ComparisonFunc::Less: + return MTL::CompareFunctionLess; + case ComparisonFunc::Equal: + return MTL::CompareFunctionEqual; + case ComparisonFunc::LessEqual: + return MTL::CompareFunctionLessEqual; + case ComparisonFunc::Greater: + return MTL::CompareFunctionGreater; + case ComparisonFunc::NotEqual: + return MTL::CompareFunctionNotEqual; + case ComparisonFunc::GreaterEqual: + return MTL::CompareFunctionGreaterEqual; + case ComparisonFunc::Always: + return MTL::CompareFunctionAlways; + default: + return MTL::CompareFunction(0); + } +} + +MTL::VertexStepFunction MetalUtil::translateVertexStepFunction(InputSlotClass slotClass) +{ + switch (slotClass) { - case Format::R32G32B32A32_TYPELESS: return PixelFormatRGBA32Float; - case Format::R32G32B32_TYPELESS: return PixelFormatInvalid; - case Format::R32G32_TYPELESS: return PixelFormatRG32Float; - case Format::R32_TYPELESS: return PixelFormatR32Float; - - case Format::R16G16B16A16_TYPELESS: return PixelFormatRGBA16Float; - case Format::R16G16_TYPELESS: return PixelFormatRG16Float; - case Format::R16_TYPELESS: return PixelFormatR16Float; - - case Format::R8G8B8A8_TYPELESS: return PixelFormatRGBA8Unorm; - case Format::R8G8_TYPELESS: return PixelFormatRG8Unorm; - case Format::R8_TYPELESS: return PixelFormatR8Unorm; - case Format::B8G8R8A8_TYPELESS: return PixelFormatBGRA8Unorm; - - case Format::R32G32B32A32_FLOAT: return PixelFormatRGBA32Float; - case Format::R32G32B32_FLOAT: return PixelFormatInvalid; - case Format::R32G32_FLOAT: return PixelFormatRG32Float; - case Format::R32_FLOAT: return PixelFormatR32Float; - - case Format::R16G16B16A16_FLOAT: return PixelFormatRGBA16Float; - case Format::R16G16_FLOAT: return PixelFormatRG16Float; - case Format::R16_FLOAT: return PixelFormatR16Float; - - case Format::R32G32B32A32_UINT: return PixelFormatRGBA32Uint; - case Format::R32G32B32_UINT: return PixelFormatInvalid; - case Format::R32G32_UINT: return PixelFormatRG32Uint; - case Format::R32_UINT: return PixelFormatR32Uint; - - case Format::R16G16B16A16_UINT: return PixelFormatRGBA16Uint; - case Format::R16G16_UINT: return PixelFormatRG16Uint; - case Format::R16_UINT: return PixelFormatR16Uint; - - case Format::R8G8B8A8_UINT: return PixelFormatRGBA8Uint; - case Format::R8G8_UINT: return PixelFormatRG8Uint; - case Format::R8_UINT: return PixelFormatR8Uint; - - case Format::R32G32B32A32_SINT: return PixelFormatRGBA32Sint; - case Format::R32G32B32_SINT: return PixelFormatInvalid; - case Format::R32G32_SINT: return PixelFormatRG32Sint; - case Format::R32_SINT: return PixelFormatR32Sint; - - case Format::R16G16B16A16_SINT: return PixelFormatRGBA16Sint; - case Format::R16G16_SINT: return PixelFormatRG16Sint; - case Format::R16_SINT: return PixelFormatR16Sint; - - case Format::R8G8B8A8_SINT: return PixelFormatRGBA8Sint; - case Format::R8G8_SINT: return PixelFormatRG8Sint; - case Format::R8_SINT: return PixelFormatR8Sint; - - case Format::R16G16B16A16_UNORM: return PixelFormatRGBA16Unorm; - case Format::R16G16_UNORM: return PixelFormatRG16Unorm; - case Format::R16_UNORM: return PixelFormatR16Unorm; - - case Format::R8G8B8A8_UNORM: return PixelFormatRGBA8Unorm; - case Format::R8G8B8A8_UNORM_SRGB: return PixelFormatRGBA8Unorm_sRGB; - case Format::R8G8_UNORM: return PixelFormatRG8Unorm; - case Format::R8_UNORM: return PixelFormatR8Unorm; - case Format::B8G8R8A8_UNORM: return PixelFormatBGRA8Unorm; - case Format::B8G8R8A8_UNORM_SRGB: return PixelFormatBGRA8Unorm_sRGB; - case Format::B8G8R8X8_UNORM: return PixelFormatInvalid; - case Format::B8G8R8X8_UNORM_SRGB: return PixelFormatInvalid; - - case Format::R16G16B16A16_SNORM: return PixelFormatRGBA16Snorm; - case Format::R16G16_SNORM: return PixelFormatRG16Snorm; - case Format::R16_SNORM: return PixelFormatR16Snorm; - - case Format::R8G8B8A8_SNORM: return PixelFormatRGBA8Snorm; - case Format::R8G8_SNORM: return PixelFormatRG8Snorm; - case Format::R8_SNORM: return PixelFormatR8Snorm; - - case Format::D32_FLOAT: return PixelFormatDepth32Float; - case Format::D16_UNORM: return PixelFormatDepth16Unorm; - case Format::D32_FLOAT_S8_UINT: return PixelFormatDepth32Float_Stencil8; - case Format::R32_FLOAT_X32_TYPELESS: return PixelFormatInvalid; - - case Format::B4G4R4A4_UNORM: return PixelFormatABGR4Unorm; - case Format::B5G6R5_UNORM: return PixelFormatB5G6R5Unorm; - case Format::B5G5R5A1_UNORM: return PixelFormatA1BGR5Unorm; - - case Format::R9G9B9E5_SHAREDEXP: return PixelFormatRGB9E5Float; - case Format::R10G10B10A2_TYPELESS: return PixelFormatInvalid; - case Format::R10G10B10A2_UINT: return PixelFormatRGB10A2Uint; - case Format::R10G10B10A2_UNORM: return PixelFormatRGB10A2Unorm; - case Format::R11G11B10_FLOAT: return PixelFormatRG11B10Float; - - case Format::BC1_UNORM: return PixelFormatBC1_RGBA; - case Format::BC1_UNORM_SRGB: return PixelFormatBC1_RGBA_sRGB; - case Format::BC2_UNORM: return PixelFormatBC2_RGBA; - case Format::BC2_UNORM_SRGB: return PixelFormatBC2_RGBA_sRGB; - case Format::BC3_UNORM: return PixelFormatBC3_RGBA; - case Format::BC3_UNORM_SRGB: return PixelFormatBC3_RGBA_sRGB; - case Format::BC4_UNORM: return PixelFormatBC4_RUnorm; - case Format::BC4_SNORM: return PixelFormatBC4_RSnorm; - case Format::BC5_UNORM: return PixelFormatBC5_RGUnorm; - case Format::BC5_SNORM: return PixelFormatBC5_RGSnorm; - case Format::BC6H_UF16: return PixelFormatBC6H_RGBUfloat; - case Format::BC6H_SF16: return PixelFormatBC6H_RGBFloat; - case Format::BC7_UNORM: return PixelFormatBC7_RGBAUnorm; - case Format::BC7_UNORM_SRGB: return PixelFormatBC7_RGBAUnorm_sRGB; - - default: return PixelFormatInvalid; + case InputSlotClass::PerVertex: + return MTL::VertexStepFunctionPerVertex; + case InputSlotClass::PerInstance: + return MTL::VertexStepFunctionPerInstance; + default: + return MTL::VertexStepFunctionPerVertex; } } diff --git a/tools/gfx/metal/metal-util.h b/tools/gfx/metal/metal-util.h index dada9be6b..617d82e29 100644 --- a/tools/gfx/metal/metal-util.h +++ b/tools/gfx/metal/metal-util.h @@ -10,8 +10,20 @@ namespace gfx { // Utility functions for Metal struct MetalUtil { - static MTL::PixelFormat getMetalPixelFormat(Format format); - static MTL::VertexFormat getMetalVertexFormat(Format format); + static NS::SharedPtr<NS::String> createString(const char* str, NS::StringEncoding encoding = NS::UTF8StringEncoding) + { + NS::SharedPtr<NS::String> nsString = NS::TransferPtr(NS::String::alloc()->init(str, encoding)); + return nsString; + } + + static NS::SharedPtr<NS::String> createStringView(void* bytes, size_t len, NS::StringEncoding encoding = NS::UTF8StringEncoding) + { + NS::SharedPtr<NS::String> nsString = NS::TransferPtr(NS::String::alloc()->init(bytes, len, encoding, false)); + return nsString; + } + + static MTL::PixelFormat translatePixelFormat(Format format); + static MTL::VertexFormat translateVertexFormat(Format format); static inline bool isDepthFormat(MTL::PixelFormat format) { @@ -30,5 +42,23 @@ struct MetalUtil } return false; } + + static MTL::SamplerMinMagFilter translateSamplerMinMagFilter(TextureFilteringMode mode); + static MTL::SamplerMipFilter translateSamplerMipFilter(TextureFilteringMode mode); + static MTL::SamplerAddressMode translateSamplerAddressMode(TextureAddressingMode mode); + static MTL::CompareFunction translateCompareFunction(ComparisonFunc func); + + static MTL::VertexStepFunction translateVertexStepFunction(InputSlotClass slotClass); + }; + +struct ScopedAutoreleasePool +{ + ScopedAutoreleasePool() { m_pool = NS::AutoreleasePool::alloc()->init(); } + ~ScopedAutoreleasePool() { m_pool->drain(); } + NS::AutoreleasePool* m_pool; +}; + +#define AUTORELEASEPOOL ScopedAutoreleasePool _pool_; + } // namespace gfx diff --git a/tools/gfx/metal/metal-vertex-layout.h b/tools/gfx/metal/metal-vertex-layout.h index 860d5678b..b0f989efd 100644 --- a/tools/gfx/metal/metal-vertex-layout.h +++ b/tools/gfx/metal/metal-vertex-layout.h @@ -14,8 +14,7 @@ namespace metal class InputLayoutImpl : public InputLayoutBase { public: - List<MTL::VertexDescriptor*> m_vertexDescs; - List<MTL::VertexBufferLayoutDescriptor*> m_bufferLayoutDescs; + NS::SharedPtr<MTL::VertexDescriptor> m_vertexDescriptor; }; } // namespace metal diff --git a/tools/gfx/render.cpp b/tools/gfx/render.cpp index f6e32fef8..0dd1e5bf5 100644 --- a/tools/gfx/render.cpp +++ b/tools/gfx/render.cpp @@ -20,6 +20,7 @@ Result SLANG_MCALL createCPUDevice(const IDevice::Desc* desc, IDevice** outDevic Result SLANG_MCALL getD3D11Adapters(List<AdapterInfo>& outAdapters); Result SLANG_MCALL getD3D12Adapters(List<AdapterInfo>& outAdapters); Result SLANG_MCALL getVKAdapters(List<AdapterInfo>& outAdapters); +Result SLANG_MCALL getMetalAdapters(List<AdapterInfo>& outAdapters); Result SLANG_MCALL getCUDAAdapters(List<AdapterInfo>& outAdapters); Result SLANG_MCALL reportD3DLiveObjects(); @@ -275,6 +276,14 @@ extern "C" SLANG_RETURN_ON_FAIL(getCUDAAdapters(adapters)); break; #endif +#if SLANG_APPLE_FAMILY + case DeviceType::Vulkan: + SLANG_RETURN_ON_FAIL(getVKAdapters(adapters)); + break; + case DeviceType::Metal: + SLANG_RETURN_ON_FAIL(getMetalAdapters(adapters)); + break; +#endif case DeviceType::CPU: return SLANG_E_NOT_IMPLEMENTED; default: @@ -330,21 +339,38 @@ extern "C" } break; #elif SLANG_APPLE_FAMILY - case DeviceType::Default: + case DeviceType::Vulkan: + { + return createVKDevice(desc, outDevice); + } case DeviceType::Metal: { return createMetalDevice(desc, outDevice); } - case DeviceType::Vulkan: + case DeviceType::Default: { - return createVKDevice(desc, outDevice); + IDevice::Desc newDesc = *desc; + newDesc.deviceType = DeviceType::Metal; + if (_createDevice(&newDesc, outDevice) == SLANG_OK) + return SLANG_OK; + newDesc.deviceType = DeviceType::Vulkan; + if (_createDevice(&newDesc, outDevice) == SLANG_OK) + return SLANG_OK; + return SLANG_FAIL; } #elif SLANG_LINUX_FAMILY && !defined(__CYGWIN__) - case DeviceType::Default: case DeviceType::Vulkan: { return createVKDevice(desc, outDevice); } + case DeviceType::Default: + { + IDevice::Desc newDesc = *desc; + newDesc.deviceType = DeviceType::Vulkan; + if (_createDevice(&newDesc, outDevice) == SLANG_OK) + return SLANG_OK; + return SLANG_FAIL; + } #endif case DeviceType::CUDA: { @@ -415,6 +441,8 @@ extern "C" return "OpenGL"; case gfx::DeviceType::Vulkan: return "Vulkan"; + case gfx::DeviceType::Metal: + return "Metal"; case gfx::DeviceType::CPU: return "CPU"; case gfx::DeviceType::CUDA: |