[gfx] metal backend skeleton (#4223)

* add metal-cpp submodule

* add metal-cpp cmake target

* gfx metal backend skeleton

* add premake support

* add foundation framework

* add metal-cpp include to premake

* update vs project file

---------

Co-authored-by: Simon Kallweit <skallweit@nvidia.com>
Co-authored-by: Jay Kwak <82421531+jkwak-work@users.noreply.github.com>

48 files changed, 3803 insertions, 3 deletions

diff --git a/tools/gfx/apple/cocoa-util.h b/tools/gfx/apple/cocoa-util.h
index 565783ee9..e9d29b87c 100644
--- a/tools/gfx/apple/cocoa-util.h
+++ b/tools/gfx/apple/cocoa-util.h
@@ -9,7 +9,7 @@ struct CocoaUtil {
 
     static void* createMetalLayer(void* nswindow);
     static void destroyMetalLayer(void* metalLayer);
-
+    static void* nextDrawable(void* metalLayer) ;
 };
 
 }
diff --git a/tools/gfx/apple/cocoa-util.mm b/tools/gfx/apple/cocoa-util.mm
index 29c3056a9..5c12e587a 100644
--- a/tools/gfx/apple/cocoa-util.mm
+++ b/tools/gfx/apple/cocoa-util.mm
@@ -22,10 +22,18 @@ void* CocoaUtil::createMetalLayer(void* nswindow)
     return layer;
 }
 
+void* CocoaUtil::nextDrawable(void* metalLayer)
+{
+    CAMetalLayer* layer = (CAMetalLayer*)metalLayer;
+    return [layer nextDrawable];
+}
+
 void CocoaUtil::destroyMetalLayer(void* metalLayer)
 {
     CAMetalLayer* layer = (CAMetalLayer*)metalLayer;
     [layer release];
 }
 
-}
\ No newline at end of file
+
+
+}
diff --git a/tools/gfx/gfx.slang b/tools/gfx/gfx.slang
index a6e603469..e14f5e282 100644
--- a/tools/gfx/gfx.slang
+++ b/tools/gfx/gfx.slang
@@ -47,6 +47,7 @@ public enum class DeviceType
     DirectX12,
     OpenGl,
     Vulkan,
+    Metal,
     CPU,
     CUDA,
     CountOf,
@@ -58,6 +59,7 @@ public enum class ProjectionStyle
     OpenGl,
     DirectX,
     Vulkan,
+    Metal,
     CountOf,
 };
 
@@ -67,6 +69,7 @@ public enum class BindingStyle
     DirectX,
     OpenGl,
     Vulkan,
+    Metal,
     CPU,
     CUDA,
     CountOf,
@@ -336,6 +339,7 @@ public enum class InteropHandleAPI
     FileDescriptor,           // A file descriptor.
     DeviceAddress,            // A device address.
     D3D12CpuDescriptorHandle, // A D3D12_CPU_DESCRIPTOR_HANDLE value.
+    Metal,                    // A general Metal object handle.
 };
 
 public struct InteropHandle
diff --git a/tools/gfx/metal/metal-api.cpp b/tools/gfx/metal/metal-api.cpp
new file mode 100644
index 000000000..aa13a605f
--- /dev/null
+++ b/tools/gfx/metal/metal-api.cpp
@@ -0,0 +1,8 @@
+// metal-api.cpp
+
+#define NS_PRIVATE_IMPLEMENTATION
+#define CA_PRIVATE_IMPLEMENTATION
+#define MTL_PRIVATE_IMPLEMENTATION
+#include <Foundation/Foundation.hpp>
+#include <QuartzCore/QuartzCore.hpp>
+#include <Metal/Metal.hpp>
diff --git a/tools/gfx/metal/metal-api.h b/tools/gfx/metal/metal-api.h
new file mode 100644
index 000000000..8bd37972a
--- /dev/null
+++ b/tools/gfx/metal/metal-api.h
@@ -0,0 +1,5 @@
+// metal-api.h
+#pragma once
+
+#include <Metal/Metal.hpp>
+#include <QuartzCore/QuartzCore.hpp>
diff --git a/tools/gfx/metal/metal-base.h b/tools/gfx/metal/metal-base.h
new file mode 100644
index 000000000..94577f145
--- /dev/null
+++ b/tools/gfx/metal/metal-base.h
@@ -0,0 +1,53 @@
+// metal-base.h
+// Shared header file for Metal implementation.
+#pragma once
+
+#include "../command-encoder-com-forward.h"
+#include "../mutable-shader-object.h"
+#include "../renderer-shared.h"
+#include "../transient-resource-heap-base.h"
+#include "core/slang-chunked-list.h"
+#include "metal-api.h"
+
+namespace gfx
+{
+namespace metal
+{
+
+    class DeviceImpl;
+    class InputLayoutImpl;
+    class BufferResourceImpl;
+    class FenceImpl;
+    class TextureResourceImpl;
+    class SamplerStateImpl;
+    class ResourceViewImpl;
+    class TextureResourceViewImpl;
+    class TexelBufferResourceViewImpl;
+    class PlainBufferResourceViewImpl;
+    class AccelerationStructureImpl;
+    class FramebufferLayoutImpl;
+    class RenderPassLayoutImpl;
+    class FramebufferImpl;
+    class PipelineStateImpl;
+    class RayTracingPipelineStateImpl;
+    class ShaderObjectLayoutImpl;
+    class EntryPointLayout;
+    class RootShaderObjectLayout;
+    class ShaderProgramImpl;
+    class PipelineCommandEncoder;
+    class ShaderObjectImpl;
+    class MutableShaderObjectImpl;
+    //class RootShaderObjectImpl;
+    class ShaderTableImpl;
+    class ResourceCommandEncoder;
+    class RenderCommandEncoder;
+    class ComputeCommandEncoder;
+    class RayTracingCommandEncoder;
+    class CommandBufferImpl;
+    class CommandQueueImpl;
+    class TransientResourceHeapImpl;
+    class QueryPoolImpl;
+    class SwapchainImpl;
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-buffer.cpp b/tools/gfx/metal/metal-buffer.cpp
new file mode 100644
index 000000000..de866bcf7
--- /dev/null
+++ b/tools/gfx/metal/metal-buffer.cpp
@@ -0,0 +1,57 @@
+// metal-buffer.cpp
+#include "metal-buffer.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+BufferResourceImpl::BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* renderer)
+    : Parent(desc)
+    , m_renderer(renderer)
+{
+    assert(renderer);
+}
+
+BufferResourceImpl::~BufferResourceImpl()
+{
+    if (sharedHandle.handleValue != 0)
+    {
+    }
+}
+
+DeviceAddress BufferResourceImpl::getDeviceAddress()
+{
+    return (DeviceAddress)0;
+}
+
+Result BufferResourceImpl::getNativeResourceHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result BufferResourceImpl::getSharedHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result BufferResourceImpl::map(MemoryRange* rangeToRead, void** outPointer)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result BufferResourceImpl::unmap(MemoryRange* writtenRange)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result BufferResourceImpl::setDebugName(const char* name)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-buffer.h b/tools/gfx/metal/metal-buffer.h
new file mode 100644
index 000000000..96a3538f2
--- /dev/null
+++ b/tools/gfx/metal/metal-buffer.h
@@ -0,0 +1,43 @@
+// metal-buffer.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class BufferResourceImpl : public BufferResource
+{
+public:
+    typedef BufferResource Parent;
+
+    BufferResourceImpl(const IBufferResource::Desc& desc, DeviceImpl* renderer);
+
+    ~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 getSharedHandle(InteropHandle* outHandle) 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;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL setDebugName(const char* name) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-buffer.cpp b/tools/gfx/metal/metal-command-buffer.cpp
new file mode 100644
index 000000000..08caa03dd
--- /dev/null
+++ b/tools/gfx/metal/metal-command-buffer.cpp
@@ -0,0 +1,86 @@
+// metal-command-buffer.cpp
+#include "metal-command-buffer.h"
+
+#include "metal-device.h"
+#include "metal-command-encoder.h"
+#include "metal-shader-object.h"
+#include "metal-command-queue.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+ICommandBuffer* CommandBufferImpl::getInterface(const Guid& guid)
+{
+    if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_ICommandBuffer)
+        return static_cast<ICommandBuffer*>(this);
+    return nullptr;
+}
+
+void CommandBufferImpl::comFree() { }
+
+Result CommandBufferImpl::init(DeviceImpl* renderer, TransientResourceHeapImpl* transientHeap)
+{
+    m_renderer = renderer;
+    m_commandBuffer = m_renderer->m_commandQueue->commandBuffer();
+    return SLANG_OK;
+}
+
+void CommandBufferImpl::encodeRenderCommands(
+    IRenderPassLayout* renderPass, IFramebuffer* framebuffer, IRenderCommandEncoder** outEncoder)
+{
+    if (!m_renderCommandEncoder)
+    {
+        m_renderCommandEncoder = new RenderCommandEncoder;
+        m_renderCommandEncoder->init(this);
+    }
+    m_renderCommandEncoder->beginPass(renderPass, framebuffer);
+    *outEncoder = m_renderCommandEncoder;
+}
+
+void CommandBufferImpl::encodeComputeCommands(IComputeCommandEncoder** outEncoder)
+{
+    if (!m_computeCommandEncoder)
+    {
+        m_computeCommandEncoder = new ComputeCommandEncoder;
+        m_computeCommandEncoder->init(this);
+    }
+    *outEncoder = m_computeCommandEncoder;
+}
+
+void CommandBufferImpl::encodeResourceCommands(IResourceCommandEncoder** outEncoder)
+{
+    if (!m_resourceCommandEncoder)
+    {
+        m_resourceCommandEncoder = new ResourceCommandEncoder;
+        m_resourceCommandEncoder->init(this);
+    }
+    *outEncoder = m_resourceCommandEncoder;
+}
+
+void CommandBufferImpl::encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder)
+{
+    if (!m_rayTracingCommandEncoder)
+    {
+        m_rayTracingCommandEncoder = new RayTracingCommandEncoder;
+        m_rayTracingCommandEncoder->init(this);
+    }
+    *outEncoder = m_rayTracingCommandEncoder;
+}
+
+void CommandBufferImpl::close()
+{
+    //m_commandBuffer->commit();
+}
+
+Result CommandBufferImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-buffer.h b/tools/gfx/metal/metal-command-buffer.h
new file mode 100644
index 000000000..aef616036
--- /dev/null
+++ b/tools/gfx/metal/metal-command-buffer.h
@@ -0,0 +1,62 @@
+// metal-command-buffer.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-shader-object.h"
+#include "metal-command-encoder.h"
+#include "../simple-transient-resource-heap.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class CommandBufferImpl
+    : public ICommandBuffer
+    , public ComObject
+{
+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;
+    RootShaderObjectImpl m_rootObject;
+
+    ResourceCommandEncoder* m_resourceCommandEncoder = nullptr;
+    ComputeCommandEncoder* m_computeCommandEncoder = nullptr;
+    RenderCommandEncoder* m_renderCommandEncoder = nullptr;
+    RayTracingCommandEncoder* m_rayTracingCommandEncoder = nullptr;
+
+    // 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);
+
+    void beginCommandBuffer();
+
+public:
+    virtual SLANG_NO_THROW void SLANG_MCALL encodeRenderCommands(
+        IRenderPassLayout* renderPass,
+        IFramebuffer* framebuffer,
+        IRenderCommandEncoder** outEncoder) override;
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        encodeComputeCommands(IComputeCommandEncoder** outEncoder) override;
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        encodeResourceCommands(IResourceCommandEncoder** outEncoder) override;
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override;
+    virtual SLANG_NO_THROW void SLANG_MCALL close() override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-encoder.cpp b/tools/gfx/metal/metal-command-encoder.cpp
new file mode 100644
index 000000000..1e74733d1
--- /dev/null
+++ b/tools/gfx/metal/metal-command-encoder.cpp
@@ -0,0 +1,475 @@
+// metal-command-encoder.cpp
+#include "metal-command-encoder.h"
+
+#include "metal-buffer.h"
+#include "metal-command-buffer.h"
+#include "metal-query.h"
+#include "metal-render-pass.h"
+#include "metal-resource-views.h"
+#include "metal-shader-object.h"
+#include "metal-shader-program.h"
+#include "metal-shader-table.h"
+#include "metal-texture.h"
+
+#include "metal-helper-functions.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+void PipelineCommandEncoder::init(CommandBufferImpl* commandBuffer)
+{
+    m_commandBuffer = commandBuffer;
+    m_device = commandBuffer->m_renderer;
+    m_metalCommandBuffer = m_commandBuffer->m_commandBuffer;
+}
+
+void ResourceCommandEncoder::copyBuffer(
+    IBufferResource* dst, Offset dstOffset, IBufferResource* src, Offset srcOffset, Size size)
+{
+}
+
+void ResourceCommandEncoder::uploadBufferData(
+    IBufferResource* buffer, Offset offset, Size size, void* data)
+{
+}
+
+void ResourceCommandEncoder::textureBarrier(
+    GfxCount count, ITextureResource* const* textures, ResourceState src, ResourceState dst)
+{
+}
+
+// TODO: Change size_t to Count?
+void ResourceCommandEncoder::bufferBarrier(
+    GfxCount count, IBufferResource* const* buffers, ResourceState src, ResourceState dst)
+{
+}
+
+void ResourceCommandEncoder::endEncoding()
+{
+}
+
+void ResourceCommandEncoder::writeTimestamp(IQueryPool* queryPool, GfxIndex index)
+{
+}
+
+void ResourceCommandEncoder::copyTexture(
+    ITextureResource* dst,
+    ResourceState dstState,
+    SubresourceRange dstSubresource,
+    ITextureResource::Offset3D dstOffset,
+    ITextureResource* src,
+    ResourceState srcState,
+    SubresourceRange srcSubresource,
+    ITextureResource::Offset3D srcOffset,
+    ITextureResource::Extents extent)
+{
+}
+
+void ResourceCommandEncoder::uploadTextureData(
+    ITextureResource* dst,
+    SubresourceRange subResourceRange,
+    ITextureResource::Offset3D offset,
+    ITextureResource::Extents extend,
+    ITextureResource::SubresourceData* subResourceData,
+    GfxCount subResourceDataCount)
+{
+}
+
+
+void ResourceCommandEncoder::clearResourceView(
+    IResourceView* view, ClearValue* clearValue, ClearResourceViewFlags::Enum flags)
+{
+}
+
+void ResourceCommandEncoder::resolveResource(
+    ITextureResource* source,
+    ResourceState sourceState,
+    SubresourceRange sourceRange,
+    ITextureResource* dest,
+    ResourceState destState,
+    SubresourceRange destRange)
+{
+}
+
+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)
+{
+}
+
+void ResourceCommandEncoder::beginDebugEvent(const char* name, float rgbColor[3])
+{
+}
+
+void ResourceCommandEncoder::endDebugEvent()
+{
+}
+
+void RenderCommandEncoder::beginPass(IRenderPassLayout* renderPass, IFramebuffer* framebuffer)
+{
+    FramebufferImpl* fb = static_cast<FramebufferImpl*>(framebuffer);
+    if (fb == nullptr)
+    {
+        return;
+    }
+    RenderPassLayoutImpl* renderPassLayoutImpl = static_cast<RenderPassLayoutImpl*>(renderPass);
+
+    MTL::RenderPassDescriptor* rpd = renderPassLayoutImpl->m_renderPassDesc->copy();
+
+    if (rpd->depthAttachment() && false)
+    {
+        TextureResourceViewImpl* depthView = static_cast<TextureResourceViewImpl*>(fb->depthStencilView.get());
+        rpd->depthAttachment()->setTexture(depthView->m_texture->m_texture);
+    }
+    const int colorTargetCount = fb->renderTargetViews.getCount();
+    for (int i = 0; i < colorTargetCount; ++i)
+    {
+        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;
+        }
+        rpd->colorAttachments()->object(i)->setTexture(tex);
+        rpd->colorAttachments()->object(i)->setClearColor(MTL::ClearColor(0.2, 0.4, 0.9, 1.0));
+    }
+    rpd->setRenderTargetWidth(fb->m_width);
+    rpd->setRenderTargetHeight(fb->m_height);
+
+    m_encoder = m_metalCommandBuffer->renderCommandEncoder(rpd);
+}
+
+void RenderCommandEncoder::endEncoding()
+{
+    m_encoder->endEncoding();
+}
+
+Result RenderCommandEncoder::bindPipeline(
+    IPipelineState* pipelineState, IShaderObject** outRootObject)
+{
+    m_currentPipeline = static_cast<PipelineStateImpl*>(pipelineState);
+    // Initialize the root object
+    SLANG_RETURN_ON_FAIL(m_commandBuffer->m_rootObject.init(m_commandBuffer->m_renderer,
+        m_currentPipeline->getProgram<ShaderProgramImpl>()->m_rootObjectLayout));
+    *outRootObject = &m_commandBuffer->m_rootObject;
+    //if (pPipelineState->m_renderState == nullptr) return SLANG_ERROR_INVALID_PARAMETER;
+    //m_encoder->setRenderPipelineState(pPipelineState->m_renderState);
+    return SLANG_OK;
+}
+
+Result RenderCommandEncoder::bindPipelineWithRootObject(
+    IPipelineState* pipelineState, IShaderObject* rootObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+void RenderCommandEncoder::setViewports(GfxCount count, const Viewport* viewports)
+{
+    static const int kMaxViewports = 8; // TODO: base on device caps
+    assert(count <= kMaxViewports);
+
+    m_viewports.setCount(count);
+    for (GfxIndex i = 0; i < count; ++i)
+    {
+        const auto& inViewport = viewports[i];
+        auto& metalViewport = m_viewports[i];
+        metalViewport.height = inViewport.extentY;
+        metalViewport.width = inViewport.extentX;
+        metalViewport.originX = inViewport.originX;
+        metalViewport.originY = inViewport.originY;
+        metalViewport.znear = inViewport.minZ;
+        metalViewport.zfar = inViewport.maxZ;
+    }
+    m_encoder->setViewports(m_viewports.begin(), count);
+}
+
+void RenderCommandEncoder::setScissorRects(GfxCount count, const ScissorRect* rects)
+{
+    static const int kMaxScissorRects = 9; // TODO:
+    assert(count < kMaxScissorRects);
+
+    m_scissorRects.setCount(count);
+    for (GfxIndex i = 0; i < count; ++i)
+    {
+        const auto& inRect = rects[i];
+        auto& metalRect = m_scissorRects[i];
+        metalRect.height = inRect.maxX - inRect.minX;
+        metalRect.width = inRect.maxY - inRect.minY;
+        metalRect.x = inRect.minX;
+        metalRect.y = inRect.minY;
+    }
+    m_encoder->setScissorRects(m_scissorRects.begin(), count);
+}
+
+void RenderCommandEncoder::setPrimitiveTopology(PrimitiveTopology topology)
+{
+}
+
+void RenderCommandEncoder::setVertexBuffers(
+    GfxIndex startSlot,
+    GfxCount slotCount,
+    IBufferResource* const* buffers,
+    const Offset* offsets)
+{
+    for (GfxIndex i = 0; i < GfxIndex(slotCount); i++)
+    {
+        GfxIndex slotIndex = startSlot + i;
+        BufferResourceImpl* buffer = static_cast<BufferResourceImpl*>(buffers[i]);
+        if (buffer)
+        {
+            MTL::Buffer* vertexBuffers = {buffer->m_buffer};
+            m_encoder->setVertexBuffer(buffer->m_buffer, offsets[i], slotIndex);
+            // ...
+        }
+    }
+}
+
+void RenderCommandEncoder::setIndexBuffer(
+    IBufferResource* buffer, Format indexFormat, Offset offset)
+{
+}
+
+Result RenderCommandEncoder::prepareDraw()
+{
+    // Bind render state, including JIT pipeline state object creation
+    auto pipeline = static_cast<PipelineStateImpl*>(m_currentPipeline.Ptr());
+    if (!pipeline)
+    {
+        return SLANG_FAIL;
+    }
+    // TODO: specialization, binding, ...
+    SLANG_RETURN_ON_FAIL(pipeline->ensureAPIPipelineStateCreated());
+    return SLANG_OK;
+}
+
+static Result translatePrimitiveType(gfx::PrimitiveType primType, MTL::PrimitiveType& mtlType)
+{
+    switch (primType) 
+    {
+    case PrimitiveType::Triangle:
+        mtlType = MTL::PrimitiveTypeTriangle;
+        break;
+    case PrimitiveType::Line:
+        mtlType = MTL::PrimitiveTypeLine;
+        break;
+    case PrimitiveType::Point:
+        mtlType = MTL::PrimitiveTypePoint;
+        break;
+    case PrimitiveType::Patch:
+    default:
+        return SLANG_E_INVALID_ARG;
+    }
+    return SLANG_OK;
+}
+
+Result RenderCommandEncoder::draw(GfxCount vertexCount, GfxIndex startVertex)
+{
+    SLANG_RETURN_ON_FAIL(prepareDraw());
+
+    MTL::PrimitiveType primType;
+    Result res = translatePrimitiveType(m_currentPipeline->desc.graphics.primitiveType, primType);
+    if (res != SLANG_OK)
+        return res;
+    m_encoder->drawPrimitives(primType, startVertex, vertexCount);
+    return SLANG_OK;
+}
+
+Result RenderCommandEncoder::drawIndexed(
+    GfxCount indexCount, GfxIndex startIndex, GfxIndex baseVertex)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+void RenderCommandEncoder::setStencilReference(uint32_t referenceValue)
+{
+}
+
+Result RenderCommandEncoder::drawIndirect(
+    GfxCount maxDrawCount,
+    IBufferResource* argBuffer,
+    Offset argOffset,
+    IBufferResource* countBuffer,
+    Offset countOffset)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RenderCommandEncoder::drawIndexedIndirect(
+    GfxCount maxDrawCount,
+    IBufferResource* argBuffer,
+    Offset argOffset,
+    IBufferResource* countBuffer,
+    Offset countOffset)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RenderCommandEncoder::setSamplePositions(
+    GfxCount samplesPerPixel, GfxCount pixelCount, const SamplePosition* samplePositions)
+{
+    return SLANG_E_NOT_AVAILABLE;
+}
+
+Result RenderCommandEncoder::drawInstanced(
+    GfxCount vertexCount,
+    GfxCount instanceCount,
+    GfxIndex startVertex,
+    GfxIndex startInstanceLocation)
+{
+    SLANG_RETURN_ON_FAIL(prepareDraw());
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RenderCommandEncoder::drawIndexedInstanced(
+    GfxCount indexCount,
+    GfxCount instanceCount,
+    GfxIndex startIndexLocation,
+    GfxIndex baseVertexLocation,
+    GfxIndex startInstanceLocation)
+{
+    SLANG_RETURN_ON_FAIL(prepareDraw());
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RenderCommandEncoder::drawMeshTasks(int x, int y, int z)
+{
+    SLANG_RETURN_ON_FAIL(prepareDraw());
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+void ComputeCommandEncoder::endEncoding() { }
+
+Result ComputeCommandEncoder::bindPipeline(
+    IPipelineState* pipelineState, IShaderObject** outRootObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ComputeCommandEncoder::bindPipelineWithRootObject(
+    IPipelineState* pipelineState, IShaderObject* rootObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ComputeCommandEncoder::dispatchCompute(int x, int y, int z)
+{
+    auto pipeline = static_cast<PipelineStateImpl*>(m_currentPipeline.Ptr());
+    if (!pipeline)
+    {
+        return SLANG_FAIL;
+    }
+
+    // Also create descriptor sets based on the given pipeline layout
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ComputeCommandEncoder::dispatchComputeIndirect(IBufferResource* argBuffer, Offset offset)
+{
+    SLANG_UNIMPLEMENTED_X("dispatchComputeIndirect");
+}
+
+void RayTracingCommandEncoder::_memoryBarrier(
+    int count,
+    IAccelerationStructure* const* structures,
+    AccessFlag srcAccess,
+    AccessFlag destAccess)
+{
+}
+
+void RayTracingCommandEncoder::_queryAccelerationStructureProperties(
+    GfxCount accelerationStructureCount,
+    IAccelerationStructure* const* accelerationStructures,
+    GfxCount queryCount,
+    AccelerationStructureQueryDesc* queryDescs)
+{
+}
+
+void RayTracingCommandEncoder::buildAccelerationStructure(
+    const IAccelerationStructure::BuildDesc& desc,
+    GfxCount propertyQueryCount,
+    AccelerationStructureQueryDesc* queryDescs)
+{
+}
+
+void RayTracingCommandEncoder::copyAccelerationStructure(
+    IAccelerationStructure* dest, IAccelerationStructure* src, AccelerationStructureCopyMode mode)
+{
+}
+
+void RayTracingCommandEncoder::queryAccelerationStructureProperties(
+    GfxCount accelerationStructureCount,
+    IAccelerationStructure* const* accelerationStructures,
+    GfxCount queryCount,
+    AccelerationStructureQueryDesc* queryDescs)
+{
+    _queryAccelerationStructureProperties(
+        accelerationStructureCount, accelerationStructures, queryCount, queryDescs);
+}
+
+void RayTracingCommandEncoder::serializeAccelerationStructure(
+    DeviceAddress dest, IAccelerationStructure* source)
+{
+}
+
+void RayTracingCommandEncoder::deserializeAccelerationStructure(
+    IAccelerationStructure* dest, DeviceAddress source)
+{
+}
+
+Result RayTracingCommandEncoder::bindPipeline(IPipelineState* pipeline, IShaderObject** outRootObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RayTracingCommandEncoder::bindPipelineWithRootObject(
+    IPipelineState* pipelineState, IShaderObject* rootObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RayTracingCommandEncoder::dispatchRays(
+    GfxIndex raygenShaderIndex,
+    IShaderTable* shaderTable,
+    GfxCount width,
+    GfxCount height,
+    GfxCount depth)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+void RayTracingCommandEncoder::endEncoding() { }
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-encoder.h b/tools/gfx/metal/metal-command-encoder.h
new file mode 100644
index 000000000..863018ca4
--- /dev/null
+++ b/tools/gfx/metal/metal-command-encoder.h
@@ -0,0 +1,327 @@
+// metal-command-encoder.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-pipeline-state.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class PipelineCommandEncoder : public ComObject
+{
+public:
+
+    void init(CommandBufferImpl* commandBuffer);
+    CommandBufferImpl* m_commandBuffer;
+    MTL::CommandBuffer* m_metalCommandBuffer = nullptr;
+    DeviceImpl* m_device = nullptr;
+    RefPtr<PipelineStateImpl> m_currentPipeline;
+};
+
+class ResourceCommandEncoder
+    : public IResourceCommandEncoder
+    , public PipelineCommandEncoder
+{
+public:
+    virtual void* getInterface(SlangUUID const& guid)
+    {
+        if (guid == GfxGUID::IID_IResourceCommandEncoder || guid == ISlangUnknown::getTypeGuid())
+            return this;
+        return nullptr;
+    }
+    virtual SLANG_NO_THROW SlangResult SLANG_MCALL
+        queryInterface(SlangUUID const& uuid, void** outObject) override
+    {
+        if (auto ptr = getInterface(uuid))
+        {
+            *outObject = ptr;
+            return SLANG_OK;
+        }
+        return SLANG_E_NO_INTERFACE;
+    }
+    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 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,
+        ResourceState dstState,
+        SubresourceRange dstSubresource,
+        ITextureResource::Offset3D dstOffset,
+        ITextureResource* src,
+        ResourceState srcState,
+        SubresourceRange srcSubresource,
+        ITextureResource::Offset3D srcOffset,
+        ITextureResource::Extents extent) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL uploadTextureData(
+        ITextureResource* dst,
+        SubresourceRange subResourceRange,
+        ITextureResource::Offset3D offset,
+        ITextureResource::Extents extend,
+        ITextureResource::SubresourceData* subResourceData,
+        GfxCount subResourceDataCount) override;
+
+    void _clearColorImage(TextureResourceViewImpl* viewImpl, ClearValue* clearValue);
+
+    void _clearDepthImage(
+        TextureResourceViewImpl* viewImpl,
+        ClearValue* clearValue,
+        ClearResourceViewFlags::Enum flags);
+
+    virtual SLANG_NO_THROW void SLANG_MCALL clearResourceView(
+        IResourceView* view, ClearValue* clearValue, ClearResourceViewFlags::Enum flags) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL resolveResource(
+        ITextureResource* source,
+        ResourceState sourceState,
+        SubresourceRange sourceRange,
+        ITextureResource* dest,
+        ResourceState destState,
+        SubresourceRange destRange) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL resolveQuery(
+        IQueryPool* queryPool,
+        GfxIndex index,
+        GfxCount count,
+        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;
+    virtual SLANG_NO_THROW void SLANG_MCALL endDebugEvent() override;
+};
+
+class RenderCommandEncoder
+    : public IRenderCommandEncoder
+    , public ResourceCommandEncoder
+{
+    SLANG_GFX_FORWARD_RESOURCE_COMMAND_ENCODER_IMPL(ResourceCommandEncoder)
+    virtual void* getInterface(SlangUUID const& uuid) override
+    {
+        if (uuid == GfxGUID::IID_IResourceCommandEncoder || uuid == GfxGUID::IID_IRenderCommandEncoder || uuid == ISlangUnknown::getTypeGuid())
+        {
+            return this;
+        }
+        return nullptr;
+    }
+public:
+    MTL::RenderCommandEncoder* m_encoder;
+    List<MTL::ScissorRect> m_scissorRects;
+    List<MTL::Viewport> m_viewports;
+
+public:
+    void beginPass(IRenderPassLayout* renderPass, IFramebuffer* framebuffer);
+
+    virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        bindPipeline(IPipelineState* pipelineState, IShaderObject** outRootObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL bindPipelineWithRootObject(
+        IPipelineState* pipelineState, IShaderObject* rootObject) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        setViewports(GfxCount count, const Viewport* viewports) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        setScissorRects(GfxCount count, const ScissorRect* rects) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        setPrimitiveTopology(PrimitiveTopology topology) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL setVertexBuffers(
+        GfxIndex startSlot,
+        GfxCount slotCount,
+        IBufferResource* const* buffers,
+        const Offset* offsets) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        setIndexBuffer(IBufferResource* buffer, Format indexFormat, Offset offset = 0) override;
+
+    Result prepareDraw();
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        draw(GfxCount vertexCount, GfxIndex startVertex = 0) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        drawIndexed(GfxCount indexCount, GfxIndex startIndex = 0, GfxIndex baseVertex = 0) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL setStencilReference(uint32_t referenceValue) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL drawIndirect(
+        GfxCount maxDrawCount,
+        IBufferResource* argBuffer,
+        Offset argOffset,
+        IBufferResource* countBuffer,
+        Offset countOffset) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL drawIndexedIndirect(
+        GfxCount maxDrawCount,
+        IBufferResource* argBuffer,
+        Offset argOffset,
+        IBufferResource* countBuffer,
+        Offset countOffset) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL setSamplePositions(
+        GfxCount samplesPerPixel,
+        GfxCount pixelCount,
+        const SamplePosition* samplePositions) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL drawInstanced(
+        GfxCount vertexCount,
+        GfxCount instanceCount,
+        GfxIndex startVertex,
+        GfxIndex startInstanceLocation) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL drawIndexedInstanced(
+        GfxCount indexCount,
+        GfxCount instanceCount,
+        GfxIndex startIndexLocation,
+        GfxIndex baseVertexLocation,
+        GfxIndex startInstanceLocation) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        drawMeshTasks(int x, int y, int z) override;
+};
+
+class ComputeCommandEncoder
+    : public IComputeCommandEncoder
+    , public ResourceCommandEncoder
+{
+public:
+    SLANG_GFX_FORWARD_RESOURCE_COMMAND_ENCODER_IMPL(ResourceCommandEncoder)
+    virtual void* getInterface(SlangUUID const& uuid) override
+    {
+        if (uuid == GfxGUID::IID_IResourceCommandEncoder || uuid == GfxGUID::IID_IComputeCommandEncoder || uuid == ISlangUnknown::getTypeGuid())
+        {
+            return this;
+        }
+        return nullptr;
+    }
+
+    virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        bindPipeline(IPipelineState* pipelineState, IShaderObject** outRootObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL bindPipelineWithRootObject(
+        IPipelineState* pipelineState, IShaderObject* rootObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL dispatchCompute(int x, int y, int z) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        dispatchComputeIndirect(IBufferResource* argBuffer, Offset offset) override;
+};
+
+class RayTracingCommandEncoder
+    : public IRayTracingCommandEncoder
+    , public ResourceCommandEncoder
+{
+public:
+    SLANG_GFX_FORWARD_RESOURCE_COMMAND_ENCODER_IMPL(ResourceCommandEncoder)
+    virtual void* getInterface(SlangUUID const& uuid) override
+    {
+        if (uuid == GfxGUID::IID_IResourceCommandEncoder || uuid == GfxGUID::IID_IRayTracingCommandEncoder || uuid == ISlangUnknown::getTypeGuid())
+        {
+            return this;
+        }
+        return nullptr;
+    }
+public:
+
+    void _memoryBarrier(
+        int count,
+        IAccelerationStructure* const* structures,
+        AccessFlag srcAccess,
+        AccessFlag destAccess);
+
+    void _queryAccelerationStructureProperties(
+        GfxCount accelerationStructureCount,
+        IAccelerationStructure* const* accelerationStructures,
+        GfxCount queryCount,
+        AccelerationStructureQueryDesc* queryDescs);
+
+    virtual SLANG_NO_THROW void SLANG_MCALL buildAccelerationStructure(
+        const IAccelerationStructure::BuildDesc& desc,
+        GfxCount propertyQueryCount,
+        AccelerationStructureQueryDesc* queryDescs) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL copyAccelerationStructure(
+        IAccelerationStructure* dest,
+        IAccelerationStructure* src,
+        AccelerationStructureCopyMode mode) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL queryAccelerationStructureProperties(
+        GfxCount accelerationStructureCount,
+        IAccelerationStructure* const* accelerationStructures,
+        GfxCount queryCount,
+        AccelerationStructureQueryDesc* queryDescs) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL
+        serializeAccelerationStructure(DeviceAddress dest, IAccelerationStructure* source) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL deserializeAccelerationStructure(
+        IAccelerationStructure* dest, DeviceAddress source) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        bindPipeline(IPipelineState* pipeline, IShaderObject** outRootObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL bindPipelineWithRootObject(
+        IPipelineState* pipelineState, IShaderObject* rootObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL dispatchRays(
+        GfxIndex raygenShaderIndex,
+        IShaderTable* shaderTable,
+        GfxCount width,
+        GfxCount height,
+        GfxCount depth) override;
+
+    virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-queue.cpp b/tools/gfx/metal/metal-command-queue.cpp
new file mode 100644
index 000000000..c8b36ff1e
--- /dev/null
+++ b/tools/gfx/metal/metal-command-queue.cpp
@@ -0,0 +1,75 @@
+// metal-command-queue.cpp
+#include "metal-command-queue.h"
+
+#include "metal-command-buffer.h"
+#include "metal-fence.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal 
+{
+
+ICommandQueue* CommandQueueImpl::getInterface(const Guid& guid)
+{
+    if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_ICommandQueue)
+        return static_cast<ICommandQueue*>(this);
+    return nullptr;
+}
+
+CommandQueueImpl::~CommandQueueImpl()
+{
+}
+
+void CommandQueueImpl::init(DeviceImpl* renderer)
+{
+    m_renderer = renderer;
+
+    MTL::Device* device = m_renderer->m_device;
+    m_commandQueue = device->newCommandQueue(8);
+}
+
+void CommandQueueImpl::waitOnHost()
+{
+}
+
+Result CommandQueueImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    outHandle->api = InteropHandleAPI::Metal;
+    outHandle->handleValue = reinterpret_cast<intptr_t>(m_commandQueue);
+    return SLANG_OK;
+}
+
+const CommandQueueImpl::Desc& CommandQueueImpl::getDesc() { return m_desc; }
+
+Result CommandQueueImpl::waitForFenceValuesOnDevice(
+    GfxCount fenceCount, IFence** fences, uint64_t* waitValues)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+void CommandQueueImpl::queueSubmitImpl(
+    uint32_t count, ICommandBuffer* const* commandBuffers, IFence* fence, uint64_t valueToSignal)
+{
+    for (uint32_t i = 0; i < count; ++i)
+    {
+        CommandBufferImpl* cmdBufImpl = static_cast<CommandBufferImpl*>(commandBuffers[i]);
+        cmdBufImpl->m_commandBuffer->presentDrawable(m_renderer->m_drawable);
+        cmdBufImpl->m_commandBuffer->commit();
+    }
+}
+
+void CommandQueueImpl::executeCommandBuffers(
+    GfxCount count, ICommandBuffer* const* commandBuffers, IFence* fence, uint64_t valueToSignal)
+{
+    if (count == 0 && fence == nullptr)
+    {
+        return;
+    }
+    queueSubmitImpl(count, commandBuffers, fence, valueToSignal);
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-command-queue.h b/tools/gfx/metal/metal-command-queue.h
new file mode 100644
index 000000000..298b0a09a
--- /dev/null
+++ b/tools/gfx/metal/metal-command-queue.h
@@ -0,0 +1,54 @@
+// metal-command-queue.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace  metal
+{
+
+class CommandQueueImpl
+    : public ICommandQueue
+    , public ComObject
+{
+public:
+    SLANG_COM_OBJECT_IUNKNOWN_ALL
+    ICommandQueue* getInterface(const Guid& guid);
+
+public:
+    Desc m_desc;
+    RefPtr<DeviceImpl> m_renderer;
+    MTL::CommandQueue* m_commandQueue = nullptr;
+    ~CommandQueueImpl();
+
+    void init(DeviceImpl* renderer);
+
+    virtual SLANG_NO_THROW void SLANG_MCALL waitOnHost() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+
+    virtual SLANG_NO_THROW const Desc& SLANG_MCALL getDesc() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL waitForFenceValuesOnDevice(
+        GfxCount fenceCount, IFence** fences, uint64_t* waitValues) override;
+
+    void queueSubmitImpl(
+        uint32_t count,
+        ICommandBuffer* const* commandBuffers,
+        IFence* fence,
+        uint64_t valueToSignal);
+
+    virtual SLANG_NO_THROW void SLANG_MCALL executeCommandBuffers(
+        GfxCount count,
+        ICommandBuffer* const* commandBuffers,
+        IFence* fence,
+        uint64_t valueToSignal) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-device.cpp b/tools/gfx/metal/metal-device.cpp
new file mode 100644
index 000000000..f20ca5d1f
--- /dev/null
+++ b/tools/gfx/metal/metal-device.cpp
@@ -0,0 +1,520 @@
+// metal-device.cpp
+#include "metal-device.h"
+
+#include "metal-swap-chain.h"
+#include "metal-util.h"
+#include "../resource-desc-utils.h"
+#include "metal-texture.h"
+#include "metal-render-pass.h"
+#include "metal-vertex-layout.h"
+#include "metal-shader-program.h"
+#include "metal-buffer.h"
+//#include "metal-command-queue.h"
+//#include "metal-fence.h"
+//#include "metal-query.h"
+//#include "metal-resource-views.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-pipeline-dump-layer.h"
+//#include "metal-helper-functions.h"
+
+#include "source/core/slang-platform.h"
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+static bool shouldDumpPipeline()
+{
+    StringBuilder dumpPipelineSettings;
+    PlatformUtil::getEnvironmentVariable(toSlice("SLANG_GFX_DUMP_PIPELINE"), dumpPipelineSettings);
+    return dumpPipelineSettings.produceString() == "1";
+}
+
+DeviceImpl::~DeviceImpl()
+{
+}
+
+Result DeviceImpl::getNativeDeviceHandles(InteropHandles* outHandles)
+{
+    outHandles->handles[0].handleValue = reinterpret_cast<intptr_t>(m_device);
+    outHandles->handles[0].api = InteropHandleAPI::Metal;
+    return SLANG_OK;
+}
+
+SlangResult DeviceImpl::initialize(const Desc& desc)
+{
+    // Initialize device info.
+    {
+        m_info.apiName = "Metal";
+        m_info.bindingStyle = BindingStyle::Metal;
+        m_info.projectionStyle = ProjectionStyle::Metal;
+        m_info.deviceType = DeviceType::Metal;
+        m_info.adapterName = "default";
+        static const float kIdentity[] = { 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1 };
+        ::memcpy(m_info.identityProjectionMatrix, kIdentity, sizeof(kIdentity));
+    }
+
+    m_desc = desc;
+
+    SLANG_RETURN_ON_FAIL(RendererBase::initialize(desc));
+    SlangResult initDeviceResult = SLANG_OK;
+
+    m_device = MTL::CreateSystemDefaultDevice();
+    m_commandQueue = m_device->newCommandQueue();
+
+    SLANG_RETURN_ON_FAIL(slangContext.initialize(
+        desc.slang,
+        desc.extendedDescCount,
+        desc.extendedDescs,
+        SLANG_METAL,
+        "sm_5_1",
+        makeArray(slang::PreprocessorMacroDesc{ "__METAL__", "1" }).getView()));
+
+    // TODO: expose via some other means
+    if (captureEnabled())
+    {
+        MTL::CaptureManager* captureManager = MTL::CaptureManager::sharedCaptureManager();
+        MTL::CaptureDescriptor* d = MTL::CaptureDescriptor::alloc()->init();
+        MTL::CaptureDestination captureDest = MTL::CaptureDestination::CaptureDestinationGPUTraceDocument;
+        if (!captureManager->supportsDestination(MTL::CaptureDestinationGPUTraceDocument))
+        {
+            std::cout << "Cannot capture MTL calls to document; ensure that Info.plist exists with 'MetalCaptureEnabled' set to 'true'." << std::endl;
+            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);
+        NS::Error* errorCode = NS::Error::alloc();
+        if (!captureManager->startCapture(d, &errorCode))
+        {
+            NS::String* errorString = errorCode->description();
+            std::string estr(errorString->cString(NS::UTF8StringEncoding));
+            std::cout << "Start capture failure: " << estr << std::endl;
+            exit(1);
+        }
+    }
+    return SLANG_OK;
+}
+
+//void DeviceImpl::waitForGpu() { m_deviceQueue.flushAndWait(); }
+
+
+SLANG_NO_THROW const DeviceInfo& SLANG_MCALL DeviceImpl::getDeviceInfo() const { return m_info; }
+
+Result DeviceImpl::createTransientResourceHeap(
+    const ITransientResourceHeap::Desc& desc, ITransientResourceHeap** outHeap)
+{
+    RefPtr<TransientResourceHeapImpl> result = new TransientResourceHeapImpl();
+    SLANG_RETURN_ON_FAIL(result->init(this, desc));
+    returnComPtr(outHeap, result);
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createCommandQueue(const ICommandQueue::Desc& desc, ICommandQueue** outQueue)
+{
+    if (m_queueAllocCount != 0)
+        return SLANG_FAIL;
+    RefPtr<CommandQueueImpl> result = new CommandQueueImpl;
+    result->init(this);
+    returnComPtr(outQueue, result);
+    m_queueAllocCount++;
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createSwapchain(
+    const ISwapchain::Desc& desc, WindowHandle window, ISwapchain** outSwapchain)
+{
+    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)
+{
+    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)
+{
+    RefPtr<RenderPassLayoutImpl> result = new RenderPassLayoutImpl;
+    SLANG_RETURN_ON_FAIL(result->init(this, desc));
+    returnComPtr(outRenderPassLayout, result);
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createFramebuffer(const IFramebuffer::Desc& desc, IFramebuffer** outFramebuffer)
+{
+    RefPtr<FramebufferImpl> fb = new FramebufferImpl;
+    SLANG_RETURN_ON_FAIL(fb->init(this, desc));
+    returnComPtr(outFramebuffer, fb);
+    return SLANG_OK;
+}
+
+SlangResult DeviceImpl::readTextureResource(
+    ITextureResource* texture,
+    ResourceState state,
+    ISlangBlob** outBlob,
+    Size* outRowPitch,
+    Size* outPixelSize)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+SlangResult DeviceImpl::readBufferResource(
+    IBufferResource* inBuffer, Offset offset, Size size, ISlangBlob** outBlob)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::getAccelerationStructurePrebuildInfo(
+    const IAccelerationStructure::BuildInputs& buildInputs,
+    IAccelerationStructure::PrebuildInfo* outPrebuildInfo)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createAccelerationStructure(
+    const IAccelerationStructure::CreateDesc& desc, IAccelerationStructure** outAS)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::getTextureAllocationInfo(
+    const ITextureResource::Desc& descIn, Size* outSize, Size* outAlignment)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::getTextureRowAlignment(Size* outAlignment)
+{
+    *outAlignment = 1;
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createTextureResource(
+    const ITextureResource::Desc& descIn,
+    const ITextureResource::SubresourceData* initData,
+    ITextureResource** outResource)
+{
+    TextureResource::Desc desc = fixupTextureDesc(descIn);
+
+    const MTL::PixelFormat format = MetalUtil::getMetalPixelFormat(desc.format);
+    if (format == 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
+    switch (desc.type)
+    {
+    case IResource::Type::Texture1D:
+    {
+        metalDesc->setTextureType(MTL::TextureType::TextureType1D);
+        metalDesc->setWidth(descIn.size.width);
+        break;
+    }
+    case IResource::Type::Texture2D:
+    {
+        metalDesc->setTextureType(MTL::TextureType::TextureType2D);
+        metalDesc->setWidth(descIn.size.width);
+        metalDesc->setHeight(descIn.size.height);
+        break;
+    }
+    case IResource::Type::TextureCube:
+    {
+        metalDesc->setTextureType(MTL::TextureType::TextureTypeCube);
+        metalDesc->setWidth(descIn.size.width);
+        metalDesc->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);
+        break;
+    }
+    default:
+    {
+        assert("!Unsupported texture type");
+        return SLANG_FAIL;
+    }
+    }
+    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);
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createBufferResource(
+    const IBufferResource::Desc& descIn, const void* initData, IBufferResource** outResource)
+{
+    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));
+
+    if (initData)
+    {
+        buffer->m_buffer = m_device->newBuffer(initData, bufferSize, opts);
+    }
+    else
+    {
+        buffer->m_buffer = m_device->newBuffer(bufferSize, opts);
+    }
+
+    returnComPtr(outResource, buffer);
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createBufferFromNativeHandle(
+    InteropHandle handle, const IBufferResource::Desc& srcDesc, IBufferResource** outResource)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createSamplerState(ISamplerState::Desc const& desc, ISamplerState** outSampler)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createTextureView(
+    ITextureResource* texture, IResourceView::Desc const& desc, IResourceView** outView)
+{
+    auto resourceImpl = static_cast<TextureResourceImpl*>(texture);
+    RefPtr<TextureResourceViewImpl> view = new TextureResourceViewImpl(this);
+    view->m_desc = desc;
+    view->m_device = this;
+    if (texture == nullptr)
+    {
+        view->m_texture = nullptr;
+        returnComPtr(outView, view);
+        return SLANG_OK;
+    }
+
+    bool isArray = resourceImpl->getDesc()->arraySize > 1;
+    MTL::PixelFormat pixelFormat = MetalUtil::getMetalPixelFormat(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;
+    switch (resourceImpl->getType())
+    {
+    case IResource::Type::Texture1D:
+        textureType = isArray ? MTL::TextureType1DArray : MTL::TextureType1D;
+        break;
+    case IResource::Type::Texture2D:
+        textureType = isArray ? MTL::TextureType2DArray : MTL::TextureType2D;
+        break;
+    case IResource::Type::Texture3D:
+    {
+        if (isArray) SLANG_UNIMPLEMENTED_X("Metal does not support arrays of 3D textures.");
+        textureType = MTL::TextureType3D;
+        break;
+    }
+    case IResource::Type::TextureCube:
+        textureType = isArray ? MTL::TextureTypeCube : MTL::TextureTypeCubeArray;
+        break;
+    default:
+        SLANG_UNIMPLEMENTED_X("Unsupported texture type.");
+        break;
+    }
+    ITextureResource::Desc newDesc = *texture->getDesc();
+    newDesc.numMipLevels = levelRange.length;
+    newDesc.arraySize = sliceRange.length;
+
+    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);
+    returnComPtr(outView, view);
+
+    return SLANG_OK;
+}
+
+Result DeviceImpl::getFormatSupportedResourceStates(Format format, ResourceStateSet* outStates)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createBufferView(
+    IBufferResource* buffer,
+    IBufferResource* counterBuffer,
+    IResourceView::Desc const& desc,
+    IResourceView** outView)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+static MTL::VertexStepFunction translateVertexStepFunction(const InputSlotClass& slotClass)
+{
+    switch (slotClass)
+    {
+    case InputSlotClass::PerInstance:       return MTL::VertexStepFunctionPerInstance;
+    case InputSlotClass::PerVertex:
+    default:                                return MTL::VertexStepFunctionPerVertex;
+    }
+}
+
+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;
+
+    const InputElementDesc* srcElements = desc.inputElements;
+    Int numElements = desc.inputElementCount;
+
+    const VertexStreamDesc* srcVertexStreams = desc.vertexStreams;
+    Int vertexStreamCount = desc.vertexStreamCount;
+
+    dstAttributes.setCount(numElements);
+    dstBufferLayouts.setCount(vertexStreamCount);
+
+    for (Int i = 0; i < vertexStreamCount; ++i)
+    {
+        auto& vbld = dstBufferLayouts[i];
+        auto& srcStream = srcVertexStreams[i];
+        vbld->setStepFunction(translateVertexStepFunction(srcStream.slotClass));
+        vbld->setStepRate(srcStream.instanceDataStepRate);
+        vbld->setStride(srcStream.stride);
+    }
+
+    for (Int i = 0; i < numElements; ++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);
+        if (metalFormat == MTL::VertexFormatInvalid)
+        {
+            return SLANG_FAIL;
+        }
+        dstAttrib->attributes()->object(i)->setFormat(metalFormat);
+    }
+
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createProgram(
+    const IShaderProgram::Desc& desc, IShaderProgram** outProgram, ISlangBlob** outDiagnosticBlob)
+{
+    // TODO:
+    RefPtr<ShaderProgramImpl> shaderProgram = new ShaderProgramImpl(this);
+    shaderProgram->init(desc);
+
+    //m_deviceObjectsWithPotentialBackReferences.add(shaderProgram);
+
+    RootShaderObjectLayout::create(
+        this,
+        shaderProgram->linkedProgram,
+        shaderProgram->linkedProgram->getLayout(),
+        shaderProgram->m_rootObjectLayout.writeRef());
+    returnComPtr(outProgram, shaderProgram);
+
+    return SLANG_OK;
+}
+
+Result DeviceImpl::createShaderObjectLayout(
+    slang::ISession* session,
+    slang::TypeLayoutReflection* typeLayout,
+    ShaderObjectLayoutBase** outLayout)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createShaderObject(ShaderObjectLayoutBase* layout, IShaderObject** outObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createMutableShaderObject(
+    ShaderObjectLayoutBase* layout, IShaderObject** outObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createMutableRootShaderObject(IShaderProgram* program, IShaderObject** outObject)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createShaderTable(const IShaderTable::Desc& desc, IShaderTable** outShaderTable)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createGraphicsPipelineState(
+    const GraphicsPipelineStateDesc& inDesc, IPipelineState** outState)
+{
+    GraphicsPipelineStateDesc desc = inDesc;
+    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)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createRayTracingPipelineState(
+    const RayTracingPipelineStateDesc& desc, IPipelineState** outState)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outPool)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::createFence(const IFence::Desc& desc, IFence** outFence)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result DeviceImpl::waitForFences(
+    GfxCount fenceCount, IFence** fences, uint64_t* fenceValues, bool waitForAll, uint64_t timeout)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-device.h b/tools/gfx/metal/metal-device.h
new file mode 100644
index 000000000..4e536ddc4
--- /dev/null
+++ b/tools/gfx/metal/metal-device.h
@@ -0,0 +1,164 @@
+// metal-device.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+#include "metal-command-buffer.h"
+#include "../simple-transient-resource-heap.h"
+#include "metal-framebuffer.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal 
+{
+
+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;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createTransientResourceHeap(
+        const ITransientResourceHeap::Desc& desc, ITransientResourceHeap** outHeap) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createCommandQueue(const ICommandQueue::Desc& desc, ICommandQueue** outQueue) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createSwapchain(
+        const ISwapchain::Desc& desc, WindowHandle window, ISwapchain** outSwapchain) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createFramebufferLayout(
+        const IFramebufferLayout::Desc& desc, IFramebufferLayout** outLayout) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createFramebuffer(const IFramebuffer::Desc& desc, IFramebuffer** outFramebuffer) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createRenderPassLayout(
+        const IRenderPassLayout::Desc& desc, IRenderPassLayout** outRenderPassLayout) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createTextureResource(
+        const ITextureResource::Desc& desc,
+        const ITextureResource::SubresourceData* initData,
+        ITextureResource** outResource) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createBufferResource(
+        const IBufferResource::Desc& desc,
+        const void* initData,
+        IBufferResource** outResource) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createBufferFromNativeHandle(
+        InteropHandle handle,
+        const IBufferResource::Desc& srcDesc,
+        IBufferResource** outResource) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createSamplerState(ISamplerState::Desc const& desc, ISamplerState** outSampler) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL createTextureView(
+        ITextureResource* texture,
+        IResourceView::Desc const& desc,
+        IResourceView** outView) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createBufferView(
+        IBufferResource* buffer,
+        IBufferResource* counterBuffer,
+        IResourceView::Desc const& desc,
+        IResourceView** outView) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createInputLayout(IInputLayout::Desc const& desc, IInputLayout** outLayout) override;
+
+    virtual Result createShaderObjectLayout(
+        slang::ISession* session,
+        slang::TypeLayoutReflection* typeLayout,
+        ShaderObjectLayoutBase** outLayout) override;
+    virtual Result createShaderObject(
+        ShaderObjectLayoutBase* layout, IShaderObject** outObject) override;
+    virtual Result createMutableShaderObject(
+        ShaderObjectLayoutBase* layout, IShaderObject** outObject) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createMutableRootShaderObject(IShaderProgram* program, IShaderObject** outObject) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createShaderTable(const IShaderTable::Desc& desc, IShaderTable** outShaderTable) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createProgram(
+        const IShaderProgram::Desc& desc,
+        IShaderProgram** outProgram,
+        ISlangBlob** outDiagnosticBlob) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createGraphicsPipelineState(
+        const GraphicsPipelineStateDesc& desc, IPipelineState** outState) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createComputePipelineState(
+        const ComputePipelineStateDesc& desc, IPipelineState** outState) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL createRayTracingPipelineState(
+        const RayTracingPipelineStateDesc& desc, IPipelineState** outState) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outPool) override;
+
+    virtual SLANG_NO_THROW SlangResult SLANG_MCALL readTextureResource(
+        ITextureResource* texture,
+        ResourceState state,
+        ISlangBlob** outBlob,
+        Size* outRowPitch,
+        Size* outPixelSize) override;
+
+    virtual SLANG_NO_THROW SlangResult SLANG_MCALL readBufferResource(
+        IBufferResource* buffer, Offset offset, Size size, ISlangBlob** outBlob) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getAccelerationStructurePrebuildInfo(
+        const IAccelerationStructure::BuildInputs& buildInputs,
+        IAccelerationStructure::PrebuildInfo* outPrebuildInfo) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL createAccelerationStructure(
+        const IAccelerationStructure::CreateDesc& desc, IAccelerationStructure** outView) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getTextureAllocationInfo(
+        const ITextureResource::Desc& desc, Size* outSize, Size* outAlignment) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getTextureRowAlignment(Size* outAlignment) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        createFence(const IFence::Desc& desc, IFence** outFence) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL waitForFences(
+        GfxCount fenceCount,
+        IFence** fences,
+        uint64_t* fenceValues,
+        bool waitForAll,
+        uint64_t timeout) override;
+
+    //void waitForGpu();
+    virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        getNativeDeviceHandles(InteropHandles* outHandles) override;
+    ~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;
+
+    //DescriptorSetAllocator descriptorSetAllocator;
+
+    uint32_t m_queueAllocCount;
+
+    // A list to hold objects that may have a strong back reference to the device
+    // instance. Because of the pipeline cache in `RendererBase`, there could be a reference
+    // cycle among `DeviceImpl`->`PipelineStateImpl`->`ShaderProgramImpl`->`DeviceImpl`.
+    // Depending on whether a `PipelineState` objects gets stored in pipeline cache, there
+    // may or may not be such a reference cycle.
+    // We need to hold strong references to any objects that may become part of the reference
+    // cycle here, so that when objects like `ShaderProgramImpl` lost all public refernces, we
+    // can always safely break the strong reference in `ShaderProgramImpl::m_device` without
+    // worrying the `ShaderProgramImpl` object getting destroyed after the completion of
+    // `DeviceImpl::~DeviceImpl()'.
+    ChunkedList<RefPtr<RefObject>, 1024> m_deviceObjectsWithPotentialBackReferences;
+
+    //RefPtr<FramebufferImpl> m_emptyFramebuffer;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-fence.cpp b/tools/gfx/metal/metal-fence.cpp
new file mode 100644
index 000000000..3c0e8edef
--- /dev/null
+++ b/tools/gfx/metal/metal-fence.cpp
@@ -0,0 +1,48 @@
+// metal-gfence.cpp
+#include "metal-fence.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+FenceImpl::FenceImpl(DeviceImpl* device)
+    : m_device(device)
+{}
+
+FenceImpl::~FenceImpl()
+{
+}
+
+Result FenceImpl::init(const IFence::Desc& desc)
+{
+    return SLANG_FAIL;
+}
+
+Result FenceImpl::getCurrentValue(uint64_t* outValue)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result FenceImpl::setCurrentValue(uint64_t value)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result FenceImpl::getSharedHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result FenceImpl::getNativeHandle(InteropHandle* outNativeHandle)
+{
+    outNativeHandle->handleValue = 0;
+    return SLANG_FAIL;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-fence.h b/tools/gfx/metal/metal-fence.h
new file mode 100644
index 000000000..af5cb7806
--- /dev/null
+++ b/tools/gfx/metal/metal-fence.h
@@ -0,0 +1,36 @@
+// metal-fence.h
+#pragma once
+
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class FenceImpl : public FenceBase
+{
+public:
+    RefPtr<DeviceImpl> m_device;
+
+    FenceImpl(DeviceImpl* device);
+
+    ~FenceImpl();
+
+    Result init(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;
+};
+
+} // namespace metal 
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-framebuffer.cpp b/tools/gfx/metal/metal-framebuffer.cpp
new file mode 100644
index 000000000..c676f44eb
--- /dev/null
+++ b/tools/gfx/metal/metal-framebuffer.cpp
@@ -0,0 +1,75 @@
+// metal-framebuffer.cpp
+#include "metal-framebuffer.h"
+#include "metal-device.h"
+#include "metal-resource-views.h"
+#include "metal-helper-functions.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+FramebufferLayoutImpl::~FramebufferLayoutImpl()
+{
+    //m_renderPass->release();
+}
+
+Result FramebufferLayoutImpl::init(DeviceImpl* renderer, 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_desc = desc;
+    return SLANG_OK;
+}
+
+FramebufferImpl::~FramebufferImpl()
+{
+}
+
+Result FramebufferImpl::init(DeviceImpl* renderer, const IFramebuffer::Desc& desc)
+{
+    m_renderer = renderer;
+    m_layout = static_cast<FramebufferLayoutImpl*>(desc.layout);
+    m_width = m_height = 1;
+
+    TextureResourceViewImpl* dsv = static_cast<TextureResourceViewImpl*>(desc.depthStencilView);
+
+    // Get frame dimensions from attachments.
+    if (dsv)
+    {
+        // If we have a depth attachment, get frame size from there.
+        auto size = dsv->m_texture->getDesc()->size;
+        auto viewDesc = dsv->getViewDesc();
+        m_width = Math::Max(1u, uint32_t(size.width >> viewDesc->subresourceRange.mipLevel));
+        m_height = Math::Max(1u, uint32_t(size.height >> viewDesc->subresourceRange.mipLevel));
+    }
+    else if (desc.renderTargetCount > 0)
+    {
+        // If we don't have a depth attachment, then we must have at least
+        // one color attachment. Get frame dimension from there.
+        auto viewImpl = static_cast<TextureResourceViewImpl*>(desc.renderTargetViews[0]);
+        auto resourceDesc = viewImpl->m_texture->getDesc();
+        auto viewDesc = viewImpl->getViewDesc();
+        auto size = resourceDesc->size;
+        m_width = Math::Max(1u, uint32_t(size.width >> viewDesc->subresourceRange.mipLevel));
+        m_height = Math::Max(1u, uint32_t(size.height >> viewDesc->subresourceRange.mipLevel));
+    }
+
+    // Initialize depthstencil and render target views
+    depthStencilView = desc.depthStencilView;
+
+    renderTargetViews.setCount(desc.renderTargetCount);
+    for (int i = 0; i < desc.renderTargetCount; ++i)
+    {
+        renderTargetViews[i] = desc.renderTargetViews[i];
+    }
+
+    return SLANG_OK;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-framebuffer.h b/tools/gfx/metal/metal-framebuffer.h
new file mode 100644
index 000000000..10ee637b8
--- /dev/null
+++ b/tools/gfx/metal/metal-framebuffer.h
@@ -0,0 +1,62 @@
+// metal-framebuffer.h
+#pragma once
+
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+enum
+{
+    kMaxRenderTargets = 8,
+    kMaxTargets = kMaxRenderTargets + 1,
+};
+
+class FramebufferLayoutImpl : public FramebufferLayoutBase
+{
+public:
+#if 0
+    MTL::RenderPassDescriptor* m_renderPass = nullptr;
+    Array<MTL::RenderPassColorAttachmentDescriptor*, kMaxTargets> m_targetDescs;
+    MTL::RenderPassDepthAttachmentDescriptor* m_depthAttachmentDesc = nullptr;
+    MTL::RenderPassStencilAttachmentDescriptor* m_stencilAttachmentDesc = nullptr;
+    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);
+};
+
+class FramebufferImpl : public FramebufferBase
+{
+public:
+    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
+    Array<MTL::RenderPassColorAttachmentDescriptor*, kMaxTargets> m_colorTargetDescs;
+    MTL::RenderPassDepthAttachmentDescriptor* m_depthAttachmentDesc = nullptr;
+    MTL::RenderPassStencilAttachmentDescriptor* m_stencilAttachmentDesc = nullptr;
+#endif
+
+public:
+    ~FramebufferImpl();
+
+    Result init(DeviceImpl* renderer, const IFramebuffer::Desc& desc);
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-helper-functions.cpp b/tools/gfx/metal/metal-helper-functions.cpp
new file mode 100644
index 000000000..e884f2e76
--- /dev/null
+++ b/tools/gfx/metal/metal-helper-functions.cpp
@@ -0,0 +1,18 @@
+// metal-helper-functions.cpp
+#include "metal-helper-functions.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+Result SLANG_MCALL createMetalDevice(const IDevice::Desc* desc, IDevice** outRenderer)
+{
+    RefPtr<metal::DeviceImpl> result = new metal::DeviceImpl();
+    SLANG_RETURN_ON_FAIL(result->initialize(*desc));
+    returnComPtr(outRenderer, result);
+    return SLANG_OK;
+}
+
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-helper-functions.h b/tools/gfx/metal/metal-helper-functions.h
new file mode 100644
index 000000000..94b005321
--- /dev/null
+++ b/tools/gfx/metal/metal-helper-functions.h
@@ -0,0 +1,13 @@
+// metal-helper-functions.h
+#pragma once
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+//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
new file mode 100644
index 000000000..aca10c12a
--- /dev/null
+++ b/tools/gfx/metal/metal-pipeline-state.cpp
@@ -0,0 +1,151 @@
+// metal-pipeline-state.cpp
+#include "metal-pipeline-state.h"
+
+#include "metal-device.h"
+#include "metal-shader-program.h"
+#include "metal-shader-object-layout.h"
+#include "metal-vertex-layout.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+PipelineStateImpl::PipelineStateImpl(DeviceImpl* 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)
+{
+    PipelineStateDesc pipelineDesc;
+    pipelineDesc.type = PipelineType::Graphics;
+    pipelineDesc.graphics = inDesc;
+    initializeBase(pipelineDesc);
+}
+
+void PipelineStateImpl::init(const ComputePipelineStateDesc& inDesc)
+{
+    PipelineStateDesc pipelineDesc;
+    pipelineDesc.type = PipelineType::Compute;
+    pipelineDesc.compute = inDesc;
+    initializeBase(pipelineDesc);
+}
+
+void PipelineStateImpl::init(const RayTracingPipelineStateDesc& inDesc)
+{
+    PipelineStateDesc pipelineDesc;
+    pipelineDesc.type = PipelineType::RayTracing;
+    pipelineDesc.rayTracing.set(inDesc);
+    initializeBase(pipelineDesc);
+}
+
+Result PipelineStateImpl::createMetalRenderPipelineState()
+{
+    MTL::RenderPipelineDescriptor* pd = MTL::RenderPipelineDescriptor::alloc()->init();
+    auto programImpl = static_cast<ShaderProgramImpl*>(m_program.Ptr());
+    if (programImpl)
+    {
+        SLANG_RETURN_ON_FAIL(programImpl->compileShaders(m_device));
+    }
+
+    const auto& programReflection = m_program->linkedProgram->getLayout();
+    const auto& composedProgram = m_program->linkedProgram;
+    for (SlangUInt i = 0; i < programReflection->getEntryPointCount(); ++i)
+    {
+        SlangStage stage = programReflection->getEntryPointByIndex(i)->getStage();
+        if (stage == SLANG_STAGE_VERTEX)
+        {
+            ComPtr<slang::IBlob> metalCode;
+            {
+                ComPtr<slang::IBlob> diagnosticsBlob;
+                SlangResult result = composedProgram->getEntryPointCode(i, 0, metalCode.writeRef(), diagnosticsBlob.writeRef());
+                if (diagnosticsBlob)
+                {
+                    std::cout << diagnosticsBlob->getBufferPointer() << std::endl;
+                }
+                //MTL::Function* f = ...
+                //RETURN_ON_FAIL(result);
+                //pd->setVertexFunction();
+            }
+        }
+        //pd->setFragmentFunction();
+    }
+    // pd->colorAttachments()->object(0)->setPixelFormat(...);
+    // pd->setDepthAttachmentPixelFormat(...);
+    // Set deftault viewport and scissor
+    // Set default rasterization state
+    // Set default framebuffer layout
+    NS::Error* error;
+    m_renderState = m_device->m_device->newRenderPipelineState(pd, &error);
+    if (m_renderState == nullptr)
+    {
+        std::cout << error->localizedDescription()->utf8String() << std::endl;
+        return SLANG_E_INVALID_ARG;
+    }
+    return SLANG_OK;
+}
+
+Result PipelineStateImpl::createMetalComputePipelineState()
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result PipelineStateImpl::ensureAPIPipelineStateCreated()
+{
+    if (m_renderState)
+        return SLANG_OK;
+
+    switch (desc.type)
+    {
+    case PipelineType::Compute:
+        return createMetalComputePipelineState();
+    case PipelineType::Graphics:
+        return createMetalRenderPipelineState();
+    default:
+        SLANG_UNREACHABLE("Unknown pipeline type.");
+        return SLANG_FAIL;
+    }
+    return SLANG_OK;
+}
+
+SLANG_NO_THROW Result SLANG_MCALL PipelineStateImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+RayTracingPipelineStateImpl::RayTracingPipelineStateImpl(DeviceImpl* device)
+    : PipelineStateImpl(device)
+{}
+
+Result RayTracingPipelineStateImpl::ensureAPIPipelineStateCreated()
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RayTracingPipelineStateImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-pipeline-state.h b/tools/gfx/metal/metal-pipeline-state.h
new file mode 100644
index 000000000..321c110a1
--- /dev/null
+++ b/tools/gfx/metal/metal-pipeline-state.h
@@ -0,0 +1,58 @@
+// metal-pipeline-state.h
+#pragma once
+
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class PipelineStateImpl : public PipelineStateBase
+{
+public:
+    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);
+
+    Result createMetalComputePipelineState();
+    Result createMetalRenderPipelineState();
+
+    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
+{
+public:
+    Dictionary<String, Index> shaderGroupNameToIndex;
+    Int shaderGroupCount;
+
+    RayTracingPipelineStateImpl(DeviceImpl* device);
+
+    virtual Result ensureAPIPipelineStateCreated() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-query.cpp b/tools/gfx/metal/metal-query.cpp
new file mode 100644
index 000000000..bfed117c6
--- /dev/null
+++ b/tools/gfx/metal/metal-query.cpp
@@ -0,0 +1,28 @@
+// metal-query.cpp
+#include "metal-query.h"
+
+//#include "metal-util.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+Result QueryPoolImpl::init(const IQueryPool::Desc& desc, DeviceImpl* device)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+QueryPoolImpl::~QueryPoolImpl()
+{
+}
+
+Result QueryPoolImpl::getResult(GfxIndex index, GfxCount count, uint64_t* data)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-query.h b/tools/gfx/metal/metal-query.h
new file mode 100644
index 000000000..02c8d3172
--- /dev/null
+++ b/tools/gfx/metal/metal-query.h
@@ -0,0 +1,30 @@
+// metal-query.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class QueryPoolImpl : public QueryPoolBase
+{
+public:
+    Result init(const IQueryPool::Desc& desc, DeviceImpl* device);
+    ~QueryPoolImpl();
+
+public:
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        getResult(GfxIndex index, GfxCount count, uint64_t* data) override;
+
+public:
+    RefPtr<DeviceImpl> m_device;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-render-pass.cpp b/tools/gfx/metal/metal-render-pass.cpp
new file mode 100644
index 000000000..b7016ebc2
--- /dev/null
+++ b/tools/gfx/metal/metal-render-pass.cpp
@@ -0,0 +1,85 @@
+// metal-render-pass.cpp
+#include "metal-render-pass.h"
+
+//#include "metal-helper-functions.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+IRenderPassLayout* RenderPassLayoutImpl::getInterface(const Guid& guid)
+{
+    if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IRenderPassLayout)
+        return static_cast<IRenderPassLayout*>(this);
+    return nullptr;
+}
+
+RenderPassLayoutImpl::~RenderPassLayoutImpl()
+{
+}
+
+static inline MTL::LoadAction translateLoadOp(IRenderPassLayout::TargetLoadOp loadOp)
+{
+    switch (loadOp)
+    {
+    case IRenderPassLayout::TargetLoadOp::Clear:
+        return MTL::LoadAction::LoadActionClear;
+    case IRenderPassLayout::TargetLoadOp::Load:
+        return MTL::LoadAction::LoadActionLoad;
+    case IRenderPassLayout::TargetLoadOp::DontCare:
+    default:
+        return MTL::LoadAction::LoadActionDontCare;
+    }
+}
+
+static inline MTL::StoreAction translateStoreOp(IRenderPassLayout::TargetStoreOp storeOp)
+{
+    switch (storeOp)
+    {
+    case IRenderPassLayout::TargetStoreOp::Store:
+        return MTL::StoreAction::StoreActionStore;
+    case IRenderPassLayout::TargetStoreOp::DontCare:
+    default:
+        return MTL::StoreAction::StoreActionDontCare;
+    }
+}
+
+Result RenderPassLayoutImpl::init(DeviceImpl* renderer, const IRenderPassLayout::Desc& desc)
+{
+    m_renderer = renderer;
+
+    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->setRenderTargetArrayLength(desc.renderTargetCount);
+
+    MTL::RenderPassColorAttachmentDescriptorArray* colorAttachments = m_renderPassDesc->colorAttachments();
+    for (GfxIndex i = 0; i < desc.renderTargetCount; ++i)
+    {
+        MTL::RenderPassColorAttachmentDescriptor* colorAttach = MTL::RenderPassColorAttachmentDescriptor::alloc()->init();
+        colorAttach->setLoadAction(translateLoadOp(desc.renderTargetAccess[i].loadOp));
+        colorAttach->setStoreAction(translateStoreOp(desc.renderTargetAccess[i].storeOp));
+        // We set the texture when the render pass is executed, using the associated framebuffer.
+        colorAttach->setTexture(nullptr);
+        colorAttachments->setObject(colorAttach, i);
+    }
+    m_renderPassDesc->depthAttachment()->setLoadAction(translateLoadOp(desc.depthStencilAccess->loadOp));
+    m_renderPassDesc->depthAttachment()->setStoreAction(translateStoreOp(desc.depthStencilAccess->storeOp));
+    // We set the depth texture when the render pass is executed, using the associated framebuffer.
+    m_renderPassDesc->depthAttachment()->setTexture(nullptr);
+    //m_renderPassDesc->depthAttachment()->setClearDepth(1000000.);
+    m_renderPassDesc->stencilAttachment()->setLoadAction(translateLoadOp(desc.depthStencilAccess->loadOp));
+    m_renderPassDesc->stencilAttachment()->setStoreAction(translateStoreOp(desc.depthStencilAccess->storeOp));
+    // We set the stencil texture when the render pass is executed, using the associated framebuffer.
+    m_renderPassDesc->stencilAttachment()->setTexture(nullptr);
+    return SLANG_OK;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-render-pass.h b/tools/gfx/metal/metal-render-pass.h
new file mode 100644
index 000000000..0ccf3424b
--- /dev/null
+++ b/tools/gfx/metal/metal-render-pass.h
@@ -0,0 +1,32 @@
+// metal-render-pass.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class RenderPassLayoutImpl
+    : public IRenderPassLayout
+    , public ComObject
+{
+public:
+    SLANG_COM_OBJECT_IUNKNOWN_ALL
+    IRenderPassLayout* getInterface(const Guid& guid);
+
+public:
+    MTL::RenderPassDescriptor* m_renderPassDesc = nullptr;
+    RefPtr<DeviceImpl> m_renderer;
+    ~RenderPassLayoutImpl();
+
+    Result init(DeviceImpl* renderer, const IRenderPassLayout::Desc& desc);
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-resource-views.cpp b/tools/gfx/metal/metal-resource-views.cpp
new file mode 100644
index 000000000..8b05b66a1
--- /dev/null
+++ b/tools/gfx/metal/metal-resource-views.cpp
@@ -0,0 +1,58 @@
+// metal-resource-views.cpp
+#include "metal-resource-views.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+TextureResourceViewImpl::~TextureResourceViewImpl()
+{
+}
+
+Result TextureResourceViewImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+TexelBufferResourceViewImpl::TexelBufferResourceViewImpl(DeviceImpl* device)
+    : ResourceViewImpl(ViewType::TexelBuffer, device)
+{}
+
+TexelBufferResourceViewImpl::~TexelBufferResourceViewImpl()
+{
+}
+
+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;
+}
+
+Result AccelerationStructureImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+AccelerationStructureImpl::~AccelerationStructureImpl()
+{
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-resource-views.h b/tools/gfx/metal/metal-resource-views.h
new file mode 100644
index 000000000..735668ad1
--- /dev/null
+++ b/tools/gfx/metal/metal-resource-views.h
@@ -0,0 +1,79 @@
+// metal-resource-views.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-buffer.h"
+#include "metal-device.h"
+#include "metal-texture.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class ResourceViewImpl : public ResourceViewBase
+{
+public:
+    enum class ViewType
+    {
+        Texture,
+        TexelBuffer,
+        PlainBuffer,
+    };
+
+public:
+    ResourceViewImpl(ViewType viewType, DeviceImpl* device)
+        : m_type(viewType)
+        , m_device(device)
+    {}
+    ViewType m_type;
+    RefPtr<DeviceImpl> m_device;
+};
+
+class TextureResourceViewImpl : public ResourceViewImpl
+{
+public:
+    TextureResourceViewImpl(DeviceImpl* device)
+        : ResourceViewImpl(ViewType::Texture, device)
+    {}
+    ~TextureResourceViewImpl();
+    RefPtr<TextureResourceImpl> m_texture;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+class TexelBufferResourceViewImpl : public ResourceViewImpl
+{
+public:
+    TexelBufferResourceViewImpl(DeviceImpl* device);
+    ~TexelBufferResourceViewImpl();
+    RefPtr<BufferResourceImpl> m_buffer;
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+class PlainBufferResourceViewImpl : public ResourceViewImpl
+{
+public:
+    PlainBufferResourceViewImpl(DeviceImpl* device);
+    RefPtr<BufferResourceImpl> m_buffer;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+class AccelerationStructureImpl : public AccelerationStructureBase
+{
+public:
+    RefPtr<BufferResourceImpl> m_buffer;
+    RefPtr<DeviceImpl> m_device;
+
+public:
+    virtual SLANG_NO_THROW DeviceAddress SLANG_MCALL getDeviceAddress() override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+    ~AccelerationStructureImpl();
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-sampler.cpp b/tools/gfx/metal/metal-sampler.cpp
new file mode 100644
index 000000000..cd94def37
--- /dev/null
+++ b/tools/gfx/metal/metal-sampler.cpp
@@ -0,0 +1,26 @@
+// metal-sampler.cpp
+#include "metal-sampler.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+SamplerStateImpl::SamplerStateImpl(DeviceImpl* device)
+    : m_device(device)
+{}
+
+SamplerStateImpl::~SamplerStateImpl()
+{
+}
+
+Result SamplerStateImpl::getNativeHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-sampler.h b/tools/gfx/metal/metal-sampler.h
new file mode 100644
index 000000000..ac6b00e62
--- /dev/null
+++ b/tools/gfx/metal/metal-sampler.h
@@ -0,0 +1,25 @@
+// metal-sampler.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class SamplerStateImpl : public SamplerStateBase
+{
+public:
+    RefPtr<DeviceImpl> m_device;
+    SamplerStateImpl(DeviceImpl* device);
+    ~SamplerStateImpl();
+    virtual SLANG_NO_THROW Result SLANG_MCALL getNativeHandle(InteropHandle* outHandle) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-object-layout.cpp b/tools/gfx/metal/metal-shader-object-layout.cpp
new file mode 100644
index 000000000..8439aa1ff
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-object-layout.cpp
@@ -0,0 +1,21 @@
+// metal-shader-object-layout.cpp
+#include "metal-shader-object-layout.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+Result RootShaderObjectLayout::create(
+    DeviceImpl* renderer,
+    slang::IComponentType* program,
+    slang::ProgramLayout* programLayout,
+    RootShaderObjectLayout** outLayout)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-object-layout.h b/tools/gfx/metal/metal-shader-object-layout.h
new file mode 100644
index 000000000..9d441f624
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-object-layout.h
@@ -0,0 +1,120 @@
+// metal-shader-object-layout.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+enum
+{
+    kMaxDescriptorSets = 32,
+};
+
+class ShaderObjectLayoutImpl : public ShaderObjectLayoutBase
+{
+public:
+	struct BindingOffset
+	{
+	};
+    struct BindingRangeInfo
+    {
+        slang::BindingType bindingType;
+        Index count;
+        Index baseIndex;
+        Index subObjectIndex;
+        uint32_t bindingOffset;
+        uint32_t offset;
+        bool isSpecializable = false;
+    };
+
+    /// Offset information for a sub-object range
+    struct SubObjectRangeOffset : BindingOffset
+    {
+        SubObjectRangeOffset() {}
+
+        SubObjectRangeOffset(slang::VariableLayoutReflection* varLayout)
+    //        : BindingOffset(varLayout)
+        { }
+        /// The offset for "pending" ordinary data related to this range
+        uint32_t pendingOrdinaryData = 0;
+    };
+
+    /// Stride information for a sub-object range
+    struct SubObjectRangeStride : BindingOffset
+    {
+        SubObjectRangeStride() {}
+
+        SubObjectRangeStride(slang::TypeLayoutReflection* typeLayout)
+        {
+            if (auto pendingLayout = typeLayout->getPendingDataTypeLayout())
+            {
+                pendingOrdinaryData = (uint32_t)pendingLayout->getStride();
+            }
+        }
+
+        /// The stride for "pending" ordinary data related to this range
+        uint32_t pendingOrdinaryData = 0;
+    };
+
+    /// Information about a logical binding range as reported by Slang reflection
+    struct SubObjectRangeInfo
+    {
+        /// The index of the binding range that corresponds to this sub-object range
+        Index bindingRangeIndex;
+
+        /// The layout expected for objects bound to this range (if known)
+        RefPtr<ShaderObjectLayoutImpl> layout;
+
+        /// The offset to use when binding the first object in this range
+        SubObjectRangeOffset offset;
+
+        /// Stride between consecutive objects in this range
+        SubObjectRangeStride stride;
+    };
+
+
+    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; }
+
+    SubObjectRangeInfo const& getSubObjectRange(Index index) { return m_subObjectRanges[index]; }
+    List<SubObjectRangeInfo> const& getSubObjectRanges() { return m_subObjectRanges; }
+protected:
+    List<BindingRangeInfo> m_bindingRangeInfo;
+	List<SubObjectRangeInfo> m_subObjectRanges;
+};
+
+class EntryPointLayout : public ShaderObjectLayoutImpl
+{
+    typedef ShaderObjectLayoutImpl Super;
+
+public:
+};
+
+class RootShaderObjectLayout : public ShaderObjectLayoutImpl
+{
+    typedef ShaderObjectLayoutImpl Super;
+
+public:
+    ~RootShaderObjectLayout();
+    static Result create(
+        DeviceImpl* renderer,
+        slang::IComponentType* program,
+        slang::ProgramLayout* programLayout,
+        RootShaderObjectLayout** outLayout);
+protected:
+public:
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-object.cpp b/tools/gfx/metal/metal-shader-object.cpp
new file mode 100644
index 000000000..d50c4021a
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-object.cpp
@@ -0,0 +1,109 @@
+// metal-shader-object.cpp
+
+#include "metal-command-buffer.h"
+#include "metal-command-encoder.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+Result ShaderObjectImpl::create(
+    IDevice* device, ShaderObjectLayoutImpl* layout, ShaderObjectImpl** outShaderObject)
+{
+    auto object = RefPtr<ShaderObjectImpl>(new ShaderObjectImpl());
+    SLANG_RETURN_ON_FAIL(object->init(device, layout));
+
+    returnRefPtrMove(outShaderObject, object);
+    return SLANG_OK;
+}
+
+RendererBase* ShaderObjectImpl::getDevice() { return m_layout->getDevice(); }
+
+GfxCount ShaderObjectImpl::getEntryPointCount() { return 0; }
+
+Result ShaderObjectImpl::getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint)
+{
+    *outEntryPoint = nullptr;
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+const void* ShaderObjectImpl::getRawData() { return m_data.getBuffer(); }
+
+Size ShaderObjectImpl::getSize() { return (Size)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;
+}
+
+Result ShaderObjectImpl::setResource(ShaderOffset const& offset, IResourceView* resourceView)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ShaderObjectImpl::setSampler(ShaderOffset const& offset, ISamplerState* sampler)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ShaderObjectImpl::setCombinedTextureSampler(
+    ShaderOffset const& offset, IResourceView* textureView, ISamplerState* sampler)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result ShaderObjectImpl::init(IDevice* device, ShaderObjectLayoutImpl* layout)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+
+
+Result EntryPointShaderObject::create(
+    IDevice* device, EntryPointLayout* layout, EntryPointShaderObject** outShaderObject)
+{
+    RefPtr<EntryPointShaderObject> object = new EntryPointShaderObject();
+    SLANG_RETURN_ON_FAIL(object->init(device, layout));
+
+    returnRefPtrMove(outShaderObject, object);
+    return SLANG_OK;
+}
+
+
+Result EntryPointShaderObject::init(IDevice* device, EntryPointLayout* layout)
+{
+    //SLANG_RETURN_ON_FAIL(Super::init(device, layout));
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+
+GfxCount RootShaderObjectImpl::getEntryPointCount() { return (GfxCount)m_entryPoints.getCount(); }
+
+Result RootShaderObjectImpl::getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint)
+{
+    returnComPtr(outEntryPoint, m_entryPoints[index]);
+    return SLANG_OK;
+}
+
+Result RootShaderObjectImpl::copyFrom(IShaderObject* object, ITransientResourceHeap* transientHeap)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RootShaderObjectImpl::collectSpecializationArgs(ExtendedShaderObjectTypeList& args)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result RootShaderObjectImpl::init(IDevice* device, RootShaderObjectLayout* layout)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-object.h b/tools/gfx/metal/metal-shader-object.h
new file mode 100644
index 000000000..cfdf5be41
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-object.h
@@ -0,0 +1,121 @@
+// metal-shader-object.h
+#pragma once
+
+#include "metal-resource-views.h"
+#include "metal-sampler.h"
+#include "metal-shader-object-layout.h"
+
+namespace gfx
+{
+
+namespace metal
+{
+
+struct CombinedTextureSamplerSlot
+{
+    RefPtr<TextureResourceViewImpl> textureView;
+    RefPtr<SamplerStateImpl> sampler;
+    operator bool() { return textureView && sampler; }
+};
+
+class ShaderObjectImpl
+    : public ShaderObjectBaseImpl<ShaderObjectImpl, ShaderObjectLayoutImpl, SimpleShaderObjectData>
+{
+public:
+    static Result create(
+        IDevice* device, ShaderObjectLayoutImpl* layout, ShaderObjectImpl** outShaderObject);
+
+    RendererBase* getDevice();
+
+    virtual SLANG_NO_THROW GfxCount SLANG_MCALL getEntryPointCount() override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) override;
+
+    virtual SLANG_NO_THROW const void* SLANG_MCALL getRawData() override;
+
+    virtual SLANG_NO_THROW Size SLANG_MCALL getSize() override;
+
+    // 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 Result SLANG_MCALL
+        setResource(ShaderOffset const& offset, IResourceView* resourceView) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        setSampler(ShaderOffset const& offset, ISamplerState* sampler) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL setCombinedTextureSampler(
+        ShaderOffset const& offset, IResourceView* textureView, ISamplerState* sampler) override;
+
+protected:
+    friend class RootShaderObjectLayout;
+
+    Result init(IDevice* device, ShaderObjectLayoutImpl* layout);
+
+public:
+};
+
+class MutableShaderObjectImpl
+    : public MutableShaderObject<MutableShaderObjectImpl, ShaderObjectLayoutImpl>
+{
+public:
+};
+
+class EntryPointShaderObject : public ShaderObjectImpl
+{
+    typedef ShaderObjectImpl Super;
+
+public:
+    static Result create(
+        IDevice* device, EntryPointLayout* layout, EntryPointShaderObject** outShaderObject);
+
+    EntryPointLayout* getLayout();
+
+protected:
+    Result init(IDevice* device, EntryPointLayout* layout);
+};
+
+
+class RootShaderObjectImpl : public ShaderObjectImpl
+{
+    using Super = ShaderObjectImpl;
+
+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();
+
+    RootShaderObjectLayout* getSpecializedLayout();
+
+    List<RefPtr<EntryPointShaderObject>> const& getEntryPoints() const;
+
+    virtual GfxCount SLANG_MCALL getEntryPointCount() override;
+    virtual Result SLANG_MCALL getEntryPoint(GfxIndex index, IShaderObject** outEntryPoint) override;
+
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        copyFrom(IShaderObject* object, ITransientResourceHeap* transientHeap) 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;
+
+public:
+    Result init(IDevice* device, RootShaderObjectLayout* layout);
+    List<RefPtr<EntryPointShaderObject>> m_entryPoints;
+};
+
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-program.cpp b/tools/gfx/metal/metal-shader-program.cpp
new file mode 100644
index 000000000..05f5f9b53
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-program.cpp
@@ -0,0 +1,49 @@
+// metal-shader-program.cpp
+#include "metal-shader-program.h"
+
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+ShaderProgramImpl::ShaderProgramImpl(DeviceImpl* device)
+    : m_device(device)
+{
+}
+
+ShaderProgramImpl::~ShaderProgramImpl()
+{
+}
+
+void ShaderProgramImpl::comFree() { }
+
+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);
+    NS::Error* error;
+    MTL::Library* library = m_device->m_device->newLibrary(nsSourceString, nullptr, &error);
+    if (library == nullptr)
+    {
+        std::cout << error->localizedDescription()->utf8String() << std::endl;
+        return SLANG_E_INVALID_ARG;
+    }
+    m_entryPointNames.add(realEntryPointName);
+    m_modules.add(library);
+    return SLANG_OK;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-program.h b/tools/gfx/metal/metal-shader-program.h
new file mode 100644
index 000000000..3846328c0
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-program.h
@@ -0,0 +1,37 @@
+// metal-shader-program.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-shader-object-layout.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+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;
+
+    virtual Result createShaderModule(
+        slang::EntryPointReflection* entryPointInfo, ComPtr<ISlangBlob> kernelCode) override;
+};
+
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-table.cpp b/tools/gfx/metal/metal-shader-table.cpp
new file mode 100644
index 000000000..8561df2f4
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-table.cpp
@@ -0,0 +1,23 @@
+// metal-shader-table.cpp
+#include "metal-shader-table.h"
+
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+RefPtr<BufferResource> ShaderTableImpl::createDeviceBuffer(
+    PipelineStateBase* pipeline,
+    TransientResourceHeapBase* transientHeap,
+    IResourceCommandEncoder* encoder)
+{
+    return RefPtr<BufferResource>(0);
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-shader-table.h b/tools/gfx/metal/metal-shader-table.h
new file mode 100644
index 000000000..3895c87be
--- /dev/null
+++ b/tools/gfx/metal/metal-shader-table.h
@@ -0,0 +1,31 @@
+// metal-shader-table.h
+#pragma once
+
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class ShaderTableImpl : public ShaderTableBase
+{
+public:
+    uint32_t m_raygenTableSize;
+    uint32_t m_missTableSize;
+    uint32_t m_hitTableSize;
+    uint32_t m_callableTableSize;
+
+    DeviceImpl* m_device;
+
+    virtual RefPtr<BufferResource> createDeviceBuffer(
+        PipelineStateBase* pipeline,
+        TransientResourceHeapBase* transientHeap,
+        IResourceCommandEncoder* encoder) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-swap-chain.cpp b/tools/gfx/metal/metal-swap-chain.cpp
new file mode 100644
index 000000000..c80ee8cf8
--- /dev/null
+++ b/tools/gfx/metal/metal-swap-chain.cpp
@@ -0,0 +1,143 @@
+// metal-swap-chain.cpp
+#include "metal-swap-chain.h"
+
+#include "metal-util.h"
+#include "../apple/cocoa-util.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+ISwapchain* SwapchainImpl::getInterface(const Guid& guid)
+{
+    if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_ISwapchain)
+        return static_cast<ISwapchain*>(this);
+    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);
+}
+
+Result SwapchainImpl::init(DeviceImpl* renderer, 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_desc = desc;
+
+    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);
+
+    createSwapchainAndImages();
+
+    return SLANG_OK;
+}
+
+Result SwapchainImpl::getImage(GfxIndex index, ITextureResource** outResource)
+{
+    if (m_images.getCount() <= (Index)index)
+        return SLANG_FAIL;
+    // TODO: iff index == current
+    m_images[index]->m_isCurrentDrawable = true;
+    returnComPtr(outResource, m_images[index]);
+    return SLANG_OK;
+}
+
+Result SwapchainImpl::resize(GfxCount width, GfxCount height)
+{
+    SLANG_UNUSED(width);
+    SLANG_UNUSED(height);
+    destroySwapchainAndImages();
+    return createSwapchainAndImages();
+}
+
+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)
+    {
+        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;
+
+    return 0;
+}
+
+Result SwapchainImpl::setFullScreenMode(bool mode) { return SLANG_FAIL; }
+
+} // namespace metal 
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-swap-chain.h b/tools/gfx/metal/metal-swap-chain.h
new file mode 100644
index 000000000..c96a5fd5e
--- /dev/null
+++ b/tools/gfx/metal/metal-swap-chain.h
@@ -0,0 +1,57 @@
+// metal-swap-chain.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-command-queue.h"
+#include "metal-device.h"
+#include "metal-texture.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class SwapchainImpl
+    : public ISwapchain
+    , public ComObject
+{
+public:
+    SLANG_COM_OBJECT_IUNKNOWN_ALL
+    ISwapchain* getInterface(const Guid& guid);
+
+public:
+    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;
+    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);
+
+    virtual SLANG_NO_THROW const Desc& SLANG_MCALL getDesc() override { return m_desc; }
+    virtual SLANG_NO_THROW Result SLANG_MCALL
+        getImage(GfxIndex index, ITextureResource** outResource) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL resize(GfxCount width, GfxCount height) override;
+    virtual SLANG_NO_THROW Result SLANG_MCALL present() override;
+    virtual SLANG_NO_THROW int SLANG_MCALL acquireNextImage() override;
+    virtual SLANG_NO_THROW bool SLANG_MCALL isOccluded() override { return false; }
+    virtual SLANG_NO_THROW Result SLANG_MCALL setFullScreenMode(bool mode) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-texture.cpp b/tools/gfx/metal/metal-texture.cpp
new file mode 100644
index 000000000..092a028b0
--- /dev/null
+++ b/tools/gfx/metal/metal-texture.cpp
@@ -0,0 +1,39 @@
+// metal-texture.cpp
+#include "metal-texture.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+TextureResourceImpl::TextureResourceImpl(const Desc& desc, DeviceImpl* device)
+    : Parent(desc)
+    , m_device(device)
+{}
+
+TextureResourceImpl::~TextureResourceImpl()
+{
+}
+
+Result TextureResourceImpl::getNativeResourceHandle(InteropHandle* outHandle)
+{
+    outHandle->api = InteropHandleAPI::Metal;
+    outHandle->handleValue = reinterpret_cast<intptr_t>(m_texture);
+    return SLANG_OK;
+}
+
+Result TextureResourceImpl::getSharedHandle(InteropHandle* outHandle)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+Result TextureResourceImpl::setDebugName(const char* name)
+{
+    return SLANG_E_NOT_IMPLEMENTED;
+}
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-texture.h b/tools/gfx/metal/metal-texture.h
new file mode 100644
index 000000000..ada0f77aa
--- /dev/null
+++ b/tools/gfx/metal/metal-texture.h
@@ -0,0 +1,37 @@
+// metal-texture.h
+#pragma once
+
+#include "metal-base.h"
+#include "metal-device.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+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;
+
+    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 setDebugName(const char* name) override;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-util.cpp b/tools/gfx/metal/metal-util.cpp
new file mode 100644
index 000000000..a8c8da6b6
--- /dev/null
+++ b/tools/gfx/metal/metal-util.cpp
@@ -0,0 +1,181 @@
+// metal-util.cpp
+#include "metal-util.h"
+#include "core/slang-math.h"
+
+#include <stdlib.h>
+#include <stdio.h>
+
+namespace gfx {
+
+using namespace MTL;
+
+MTL::VertexFormat MetalUtil::getMetalVertexFormat(Format format)
+{
+    switch (format)
+    {
+    case Format::R8G8_UINT:                 return VertexFormatUChar2;
+    // VertexFormatUChar3
+    case Format::R8G8B8A8_UINT:             return VertexFormatUChar4;
+    case Format::R8G8_SINT:                 return VertexFormatChar2;
+    // return VertexFormatChar3
+    case Format::R8G8B8A8_SINT:             return VertexFormatChar4;
+    case Format::R8G8_UNORM:                return VertexFormatUChar2Normalized;
+    // return VertexFormatUChar3Normalized;
+    case Format::R8G8B8A8_UNORM:            return VertexFormatUChar4Normalized;
+    case Format::R8G8_SNORM:                return VertexFormatChar2Normalized;
+    // return VertexFormatChar3Normalized
+    case Format::R8G8B8A8_SNORM:            return VertexFormatChar4Normalized;
+    case Format::R16G16_UINT:               return VertexFormatUShort2;
+    // return VertexFormatUShort3;
+    case Format::R16G16B16A16_UINT:         return VertexFormatUShort4;
+    case Format::R16G16_SINT:               return VertexFormatShort2;
+    // return VertexFormatShort3;
+    case Format::R16G16B16A16_SINT:         return VertexFormatShort4;
+    case Format::R16G16_UNORM:              return VertexFormatUShort2Normalized;
+    // return VertexFormatUShort3Normalized;
+    case Format::R16G16B16A16_UNORM:        return VertexFormatUShort4Normalized;
+    case Format::R16G16_SNORM:              return VertexFormatShort2Normalized;
+    // return VertexFormatShort3Normalized;
+    case Format::R16G16B16A16_SNORM:        return VertexFormatShort4Normalized;
+    case Format::R16G16_FLOAT:              return 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;
+    // 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;
+    }
+}
+
+/* static */MTL::PixelFormat MetalUtil::getMetalPixelFormat(Format format)
+{
+    switch (format)
+    {
+    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;
+    }
+}
+
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-util.h b/tools/gfx/metal/metal-util.h
new file mode 100644
index 000000000..dada9be6b
--- /dev/null
+++ b/tools/gfx/metal/metal-util.h
@@ -0,0 +1,34 @@
+// metal-util.h
+#pragma once
+
+#include "core/slang-basic.h"
+#include "metal-api.h"
+#include "slang-gfx.h"
+
+namespace gfx {
+
+// Utility functions for Metal
+struct MetalUtil 
+{
+    static MTL::PixelFormat getMetalPixelFormat(Format format);
+    static MTL::VertexFormat getMetalVertexFormat(Format format);
+
+    static inline bool isDepthFormat(MTL::PixelFormat format)
+    {
+        switch (format)
+        {
+            return true;
+        }
+        return false;
+    }
+
+    static inline bool isStencilFormat(MTL::PixelFormat format)
+    {
+        switch (format)
+        {
+            return true;
+        }
+        return false;
+    }
+};
+} // namespace gfx
diff --git a/tools/gfx/metal/metal-vertex-layout.h b/tools/gfx/metal/metal-vertex-layout.h
new file mode 100644
index 000000000..860d5678b
--- /dev/null
+++ b/tools/gfx/metal/metal-vertex-layout.h
@@ -0,0 +1,22 @@
+// metal-vertex-layout.h
+#pragma once
+
+#include "metal-base.h"
+
+namespace gfx
+{
+
+using namespace Slang;
+
+namespace metal
+{
+
+class InputLayoutImpl : public InputLayoutBase
+{
+public:
+    List<MTL::VertexDescriptor*> m_vertexDescs;
+    List<MTL::VertexBufferLayoutDescriptor*> m_bufferLayoutDescs;
+};
+
+} // namespace metal
+} // namespace gfx
diff --git a/tools/gfx/render.cpp b/tools/gfx/render.cpp
index 68f81fba5..f6e32fef8 100644
--- a/tools/gfx/render.cpp
+++ b/tools/gfx/render.cpp
@@ -13,6 +13,7 @@ using namespace Slang;
 Result SLANG_MCALL createD3D11Device(const IDevice::Desc* desc, IDevice** outDevice);
 Result SLANG_MCALL createD3D12Device(const IDevice::Desc* desc, IDevice** outDevice);
 Result SLANG_MCALL createVKDevice(const IDevice::Desc* desc, IDevice** outDevice);
+Result SLANG_MCALL createMetalDevice(const IDevice::Desc* desc, IDevice** outDevice);
 Result SLANG_MCALL createCUDADevice(const IDevice::Desc* desc, IDevice** outDevice);
 Result SLANG_MCALL createCPUDevice(const IDevice::Desc* desc, IDevice** outDevice);
 
@@ -328,7 +329,17 @@ extern "C"
                 return SLANG_FAIL;
             }
             break;
-#elif (SLANG_LINUX_FAMILY || SLANG_APPLE_FAMILY) && !defined(__CYGWIN__)
+#elif SLANG_APPLE_FAMILY
+        case DeviceType::Default:
+        case DeviceType::Metal:
+        {
+            return createMetalDevice(desc, outDevice);
+        }
+        case DeviceType::Vulkan:
+        {
+            return createVKDevice(desc, outDevice);
+        }
+#elif SLANG_LINUX_FAMILY && !defined(__CYGWIN__)
         case DeviceType::Default:
         case DeviceType::Vulkan:
         {