work on gfx metal backend (#4287)

* implement sampler state

* implement input layout

* implement fence object

* buffer implementation

* texture implementation

* cleanup

* add adapter enumeration

* supported formats and allocation info

* work on device and implement readBufferResource

* skeleton for transient resource heap

* initial work on command queue / buffers / encoders

* fix uploading initial buffer data

* implement buffer resource view

* string utility functions

* wip query pool implementation

* cleanup

* swapchain

* wip

* remove plain buffer view

* extend gfxGetDeviceTypeName with metal

* basic support for resource binding with compute shaders

* needed for metal bindings

* replace assert(0) with SLANG_UNIMPLEMENTED_X

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], &registryID, 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: