diff --git a/Source/Metal/CommandAllocatorMTL.h b/Source/Metal/CommandAllocatorMTL.h
index 127c7ad..b0071a6 100644
--- a/Source/Metal/CommandAllocatorMTL.h
+++ b/Source/Metal/CommandAllocatorMTL.h
@@ -31,7 +31,7 @@ struct CommandAllocatorMTL {
 
 private:
     DeviceMTL& m_Device;
-    struct CommandQueueMTL* m_CommandQueue;
+    const struct CommandQueueMTL* m_CommandQueue;
 };
 
 }
diff --git a/Source/Metal/CommandAllocatorMTL.mm b/Source/Metal/CommandAllocatorMTL.mm
index 9fc54bd..913a523 100644
--- a/Source/Metal/CommandAllocatorMTL.mm
+++ b/Source/Metal/CommandAllocatorMTL.mm
@@ -11,7 +11,8 @@
 }
 
 Result CommandAllocatorMTL::Create(const CommandQueue& commandQueue) {
-    m_CommandQueue = &(CommandQueueMTL&)commandQueue;
+    const CommandQueueMTL& commandQueueImpl = (CommandQueueMTL&)commandQueue;
+    m_CommandQueue = &commandQueueImpl;
     return Result::SUCCESS;
 }
 
@@ -22,6 +23,7 @@
 Result CommandAllocatorMTL::CreateCommandBuffer(CommandBuffer*& commandBuffer) {
 
     CommandBufferMTL* commandBufferImpl = Allocate<CommandBufferMTL>(m_Device.GetStdAllocator(), m_Device);
+    commandBufferImpl->Create(m_CommandQueue);
     commandBuffer = (CommandBuffer*)commandBufferImpl;
     return Result::SUCCESS;
 }
diff --git a/Source/Metal/CommandBufferMTL.h b/Source/Metal/CommandBufferMTL.h
index 405dccc..18dd25b 100644
--- a/Source/Metal/CommandBufferMTL.h
+++ b/Source/Metal/CommandBufferMTL.h
@@ -14,6 +14,13 @@ struct PipelineLayoutMTL;
 struct TextureMTL;
 struct DescriptorMTL;
 
+NriBits(BarrierBits, uint8_t,
+    NONE = 0,
+    BARRIER_FLAG_BUFFERS = NriBit(0),
+    BARRIER_FLAG_TEXTURES = NriBit(1),
+    BARRIER_FLAG_RENDERTARGETS = NriBit(2),
+    BARRIER_FLAG_FENCE = NriBit(3));
+
 
 NriBits(CommandBufferDirtyBits, uint8_t,
         NONE                            = 0,
@@ -29,12 +36,11 @@ struct CommandBufferMTL {
     inline DeviceMTL& GetDevice() const {
         return m_Device;
     }
-    
-    inline operator id<MTLCommandBuffer>() const {
+ 
+    inline id<MTLCommandBuffer> GetHandle() const {
         return m_Handle;
     }
     
-    
     void SetDebugName(const char* name);
     Result Begin(const DescriptorPool* descriptorPool);
     Result End();
@@ -87,7 +93,7 @@ struct CommandBufferMTL {
     
     ~CommandBufferMTL();
     
-    void Create(id<MTLCommandBuffer> cmd);
+    void Create(const struct CommandQueueMTL* queue);
     
     struct CmdIndexBuffer {
         size_t m_Offset;
@@ -95,18 +101,31 @@ struct CommandBufferMTL {
         struct BufferMTL* m_Buffer;
     };
     
+    struct CmdVertexBuffer {
+        size_t m_Offset;
+        struct BufferMTL* m_Buffer;
+    };
+    
 private:
     
+    void Restore();
     void updateCommandBufferState();
+    void EndCurrentEncoders();
     
     DeviceMTL& m_Device;
     struct PipelineMTL* m_CurrentPipeline = nullptr;
-    struct CmdIndexBuffer m_CurrentIndexCmd;
     PipelineLayoutMTL* m_CurrentPipelineLayout = nullptr;
     id<MTLCommandBuffer> m_Handle;
     id<MTLRenderCommandEncoder> m_RendererEncoder = nil;
     id<MTLComputeCommandEncoder> m_ComputeEncoder = nil;
     id<MTLBlitCommandEncoder> m_BlitEncoder = nil;
+    const struct CommandQueueMTL* m_CommandQueue = nullptr;
+    
+    struct CmdIndexBuffer m_CurrentIndexCmd;
+    uint32_t m_dirtyVertexBufferBits = 0;
+    struct CmdVertexBuffer m_CurrentVertexCmd[32];
+    id<MTLRenderPipelineState> m_GraphicsPipelineState;
+    
     
     CommandBufferDirtyBits m_DirtyBits = CommandBufferDirtyBits::NONE;
 };
diff --git a/Source/Metal/CommandBufferMTL.mm b/Source/Metal/CommandBufferMTL.mm
index 47f17c8..13e4d9d 100644
--- a/Source/Metal/CommandBufferMTL.mm
+++ b/Source/Metal/CommandBufferMTL.mm
@@ -25,13 +25,10 @@
 }
 
 Result CommandBufferMTL::Begin(const DescriptorPool* descriptorPool) {
-    m_ComputeEncoder = [m_Handle computeCommandEncoderWithDescriptor: NULL];
 }
 
 Result CommandBufferMTL::End() {
-    m_ComputeEncoder = nil;
-    m_RendererEncoder = nil;
-    
+    EndCurrentEncoders();
 }
 
 void CommandBufferMTL::SetPipeline(const Pipeline& pipeline) {
@@ -39,6 +36,25 @@
         return;
     PipelineMTL& pipelineImpl = (PipelineMTL&)pipeline;
     m_CurrentPipeline = &pipelineImpl;
+    
+    switch(m_CurrentPipeline->GetPipelineType()) {
+        case PipelineType::Compute: {
+            if(!m_ComputeEncoder) {
+                MTLComputePassDescriptor* computePassDescriptor = [MTLComputePassDescriptor computePassDescriptor];
+                m_ComputeEncoder = [m_Handle computeCommandEncoderWithDescriptor:computePassDescriptor];
+            }
+            [m_ComputeEncoder setComputePipelineState: m_CurrentPipeline->GetComputePipeline()];
+            break;
+        }
+        case PipelineType::Graphics: {
+            [m_RendererEncoder setRenderPipelineState: m_CurrentPipeline->GetGraphicsPipeline()];
+            m_GraphicsPipelineState = m_CurrentPipeline->GetGraphicsPipeline();
+            break;
+        }
+        case PipelineType::Raytracing: {
+            break;
+        }
+    }
 }
 
 void CommandBufferMTL::SetPipelineLayout(const PipelineLayout& pipelineLayout) {
@@ -51,54 +67,57 @@
     auto* layout = m_CurrentPipelineLayout->GetDescriptorSetLayout(setIndexInPipelineLayout);
     const uint32_t space = layout->m_DescriptorSetDesc.registerSpace;
     
-    bool setFragmentBuffer = false;
-    bool setVertexBuffer = false;
-    
+    StageBits stageBits = (StageBits)0;
     for(size_t i = 0; i < layout->m_DescriptorSetDesc.rangeNum; i++) {
-        if(layout->m_DescriptorSetDesc.ranges[i].shaderStages == StageBits::ALL) {
-            setFragmentBuffer = true;
-            setVertexBuffer = true;
-            break;
-        }
-        if(layout->m_DescriptorSetDesc.ranges[i].shaderStages & StageBits::VERTEX_SHADER) {
-            setVertexBuffer = true;
-            break;
-        }
-        if(layout->m_DescriptorSetDesc.ranges[i].shaderStages & StageBits::GRAPHICS_SHADERS) {
-            setFragmentBuffer = true;
-            break;
-        }
+        if(layout->m_DescriptorSetDesc.ranges[i].shaderStages == StageBits::NONE)
+            continue;
+        stageBits |= layout->m_DescriptorSetDesc.ranges[i].shaderStages ;
     }
-    if(setFragmentBuffer) {
+    if(stageBits == StageBits::ALL ||
+       stageBits & StageBits::FRAGMENT_SHADER) {
         [m_RendererEncoder setFragmentBuffer: descriptorSetImpl.GetArgumentBuffer()
                                       offset: descriptorSetImpl.GetArugmentBufferOffset()
                                      atIndex: space];
     }
-    if(setVertexBuffer) {
+    if(stageBits == StageBits::ALL ||
+       stageBits & StageBits::VERTEX_SHADER) {
         [m_RendererEncoder setVertexBuffer: descriptorSetImpl.GetArgumentBuffer()
                                     offset: descriptorSetImpl.GetArugmentBufferOffset()
                                    atIndex: space];
         
     }
     
-
+    if(stageBits == StageBits::ALL ||
+       stageBits & StageBits::COMPUTE_SHADER) {
+        [m_ComputeEncoder setBuffer: descriptorSetImpl.GetArgumentBuffer()
+                                    offset: descriptorSetImpl.GetArugmentBufferOffset()
+                                   atIndex: space];
+        
+    }
+    
 }
-void CommandBufferMTL::SetConstants(uint32_t pushConstantIndex, const void* data, uint32_t size) {
-    //if (pDesc->mUsedStages & SHADER_STAGE_VERT)
-    //{
-    //    [m_RendererEncoder setVertexBytes:data length:size atIndex:pushConstantIndex];
-    //}
-
-    //if (pDesc->mUsedStages & SHADER_STAGE_FRAG)
-    //{
-    //    [m_RendererEncoder  setFragmentBytes:data length:size atIndex:pushConstantIndex];
-    //}
-
-    //if (pDesc->mUsedStages & SHADER_STAGE_COMP)
-    //{
-    //    [m_RendererEncoder setBytes:data length:size atIndex:pushConstantIndex];
-    //}
 
+void CommandBufferMTL::SetConstants(uint32_t pushConstantIndex, const void* data, uint32_t size) {
+    const struct RootConstantDesc* pDesc = m_CurrentPipelineLayout->GetPushBinding(pushConstantIndex);
+    
+    
+    if (pDesc->shaderStages == StageBits::ALL ||
+        pDesc->shaderStages & StageBits::VERTEX_SHADER)
+    {
+        [m_RendererEncoder setVertexBytes:data length:size atIndex:pDesc->registerIndex];
+    }
+    
+    if (pDesc->shaderStages == StageBits::ALL ||
+        pDesc->shaderStages & StageBits::FRAGMENT_SHADER )
+    {
+        [m_RendererEncoder setFragmentBytes:data length:size atIndex:pDesc->registerIndex];
+    }
+    
+    if (pDesc->shaderStages == StageBits::ALL ||
+        pDesc->shaderStages & StageBits::COMPUTE_SHADER )
+    {
+        [m_ComputeEncoder setBytes:data length:size atIndex:pDesc->registerIndex];
+    }
 }
 
 void CommandBufferMTL::SetDescriptorPool(const DescriptorPool& descriptorPool) {
@@ -107,59 +126,54 @@
 
 void CommandBufferMTL::Barrier(const BarrierGroupDesc& barrierGroupDesc) {
     
-    //if (pCmd->pQueue->mBarrierFlags & BARRIER_FLAG_BUFFERS)
-    {
-        [m_RendererEncoder memoryBarrierWithScope:MTLBarrierScopeBuffers
-                                         afterStages:MTLRenderStageFragment
-                                        beforeStages:MTLRenderStageVertex];
-    }
-
-    //if (pCmd->pQueue->mBarrierFlags & BARRIER_FLAG_TEXTURES)
-    {
-        [m_RendererEncoder memoryBarrierWithScope:MTLBarrierScopeTextures
-                                         afterStages:MTLRenderStageFragment
-                                        beforeStages:MTLRenderStageVertex];
+    const size_t totalResources = barrierGroupDesc.bufferNum + barrierGroupDesc.textureNum;
+    
+    Scratch<id<MTLResource>> resourceBarrier = AllocateScratch(m_Device, id<MTLResource>, totalResources);
+    size_t barrierCount = 0;
+    for(size_t i = 0; i < barrierGroupDesc.bufferNum; i++) {
+        const BufferBarrierDesc& in = barrierGroupDesc.buffers[i];
+        const BufferMTL& bufferImpl = *(const BufferMTL*)in.buffer;
+        resourceBarrier[barrierCount++] = bufferImpl.GetHandle();
     }
-
-    //if (pCmd->pQueue->mBarrierFlags & BARRIER_FLAG_RENDERTARGETS)
-    {
-        [m_RendererEncoder memoryBarrierWithScope:MTLBarrierScopeRenderTargets
-                                         afterStages:MTLRenderStageFragment
-                                        beforeStages:MTLRenderStageVertex];
+    
+    for(size_t i = 0; i < barrierGroupDesc.textureNum; i++) {
+        const TextureBarrierDesc& in = barrierGroupDesc.textures[i];
+        const TextureMTL& textureImpl = *(const TextureMTL*)in.texture;
+        resourceBarrier[barrierCount++] = textureImpl.GetHandle();
     }
-
     
+    [m_RendererEncoder memoryBarrierWithResources: resourceBarrier
+                                            count: barrierCount
+                                      afterStages: MTLRenderStageFragment
+                                     beforeStages: MTLRenderStageVertex];
 }
 void CommandBufferMTL::BeginRendering(const AttachmentsDesc& attachmentsDesc) {
-    MTLRenderPassDescriptor* renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
+    MTLRenderPassDescriptor* renderPassDesc = [MTLRenderPassDescriptor renderPassDescriptor];
     
     for(uint32_t i = 0; i < attachmentsDesc.colorNum; i++) {
         DescriptorMTL& descriptorMTL = *(DescriptorMTL*)attachmentsDesc.colors[i];
         
-        renderPassDescriptor.colorAttachments[i].texture = descriptorMTL.GetTextureHandle();
-        renderPassDescriptor.colorAttachments[i].clearColor = MTLClearColorMake(0, 0, 0, 1);
-        renderPassDescriptor.colorAttachments[i].loadAction = MTLLoadActionClear;
-        renderPassDescriptor.colorAttachments[i].storeAction = MTLStoreActionStore;
+        renderPassDesc.colorAttachments[i].texture = descriptorMTL.GetTextureHandle();
+        renderPassDesc.colorAttachments[i].clearColor = MTLClearColorMake(0, 0, 0, 1);
+        renderPassDesc.colorAttachments[i].loadAction = MTLLoadActionLoad;
+        renderPassDesc.colorAttachments[i].storeAction = MTLStoreActionStore;
     }
     
     if(attachmentsDesc.depthStencil) {
-        
         DescriptorMTL& descriptorMTL = *(DescriptorMTL*)attachmentsDesc.depthStencil;
-        renderPassDescriptor.depthAttachment.texture = descriptorMTL.GetTextureHandle();
-       // renderPassDescriptor.depthAttachment.clearColor = MTLClearColorMake(0, 0, 0, 1);
-        renderPassDescriptor.depthAttachment.loadAction = MTLLoadActionClear;
-        renderPassDescriptor.depthAttachment.storeAction = MTLStoreActionStore;
+        renderPassDesc.depthAttachment.texture = descriptorMTL.GetTextureHandle();
+        renderPassDesc.depthAttachment.loadAction = MTLLoadActionLoad;
+        renderPassDesc.depthAttachment.storeAction = MTLStoreActionStore;
 
     }
- 
-    m_RendererEncoder = [m_Handle renderCommandEncoderWithDescriptor: renderPassDescriptor];
+    //m_RenderPassDescriptor = [renderPassDescriptor copy];
+    m_RendererEncoder = [m_Handle renderCommandEncoderWithDescriptor: renderPassDesc];
 }
+
 void CommandBufferMTL::EndRendering() {
-    NSCAssert(m_RendererEncoder, @"Renderer Encoderer Not Set");
-    [m_RendererEncoder endEncoding];
-    m_RendererEncoder = nil;
-    m_ComputeEncoder = nil;
+    EndCurrentEncoders();
 }
+
 void CommandBufferMTL::SetViewports(const Viewport* viewports, uint32_t viewportNum) {
     Scratch<MTLViewport> mtlViewports = AllocateScratch(m_Device, MTLViewport, viewportNum);
     for(size_t i = 0; i < viewportNum; i++) {
@@ -172,6 +186,8 @@
     }
     [m_RendererEncoder setViewports: mtlViewports count: viewportNum];
 }
+
+
 void CommandBufferMTL::SetScissors(const Rect* rects, uint32_t rectNum) {
     NSCAssert(m_RendererEncoder, @"encoder set");
     MTLScissorRect rect;
@@ -197,12 +213,32 @@
     ];
 }
 void CommandBufferMTL::SetShadingRate(const ShadingRateDesc& shadingRateDesc) {
+
     //[m_RendererEncoder sha]
 }
 
 void CommandBufferMTL::ClearAttachments(const ClearDesc* clearDescs, uint32_t clearDescNum, const Rect* rects, uint32_t rectNum) {
+  
+
+    Restore();
+}
+
+void CommandBufferMTL::EndCurrentEncoders() {
     
+    if(m_RendererEncoder) {
+        [m_RendererEncoder endEncoding];
+        m_RendererEncoder = nil;
+    }
+    
+    if(m_ComputeEncoder) {
+        [m_ComputeEncoder endEncoding];
+        m_ComputeEncoder = nil;
+    }
     
+    if(m_BlitEncoder) {
+        [m_BlitEncoder endEncoding];
+        m_BlitEncoder = nil;
+    }
 }
 
 void CommandBufferMTL::SetIndexBuffer(const Buffer& buffer, uint64_t offset, IndexType indexType) {
@@ -219,9 +255,26 @@
     m_CurrentIndexCmd.m_Offset = offset;
 }
 
+
+void CommandBufferMTL::Restore() {
+    [m_RendererEncoder setRenderPipelineState: m_GraphicsPipelineState];
+    uint32_t vertexSlot = 0;
+    for( uint32_t attr = m_dirtyVertexBufferBits; attr > 0; attr = ( attr >> 1 ), vertexSlot++ ) {
+        [m_RendererEncoder setVertexBuffer: m_CurrentVertexCmd[vertexSlot].m_Buffer->GetHandle()
+                                    offset: m_CurrentVertexCmd[vertexSlot].m_Offset
+                                   atIndex: vertexSlot];
+    }
+}
+
 void CommandBufferMTL::SetVertexBuffers(uint32_t baseSlot, uint32_t bufferNum, const Buffer* const* buffers, const uint64_t* offsets) {
     for(size_t i = 0; i < bufferNum; i++) {
         BufferMTL* mtlBuffer = (BufferMTL*)buffers[i];
+        
+        const size_t slotIndex = i + baseSlot;
+        m_CurrentVertexCmd[slotIndex].m_Offset = offsets[i];
+        m_CurrentVertexCmd[slotIndex].m_Buffer = mtlBuffer;
+        m_dirtyVertexBufferBits |= (1 << slotIndex);
+        
         [m_RendererEncoder setVertexBuffer: mtlBuffer->GetHandle()
                                 offset: offsets[i]
                                atIndex: i + baseSlot];
@@ -230,16 +283,20 @@
 }
 
 void CommandBufferMTL::Draw(const DrawDesc& drawDesc) {
-    [m_RendererEncoder drawPrimitives: m_CurrentPipeline->m_primitiveType
+    //m_RendererEncoder = [m_Handle renderCommandEncoderWithDescriptor: m_renderPassDescriptor];
+     
+    [m_RendererEncoder drawPrimitives: m_CurrentPipeline->GetPrimitiveType()
             vertexStart:drawDesc.baseVertex
             vertexCount:drawDesc.vertexNum
             instanceCount:drawDesc.instanceNum
             baseInstance: 0];
+    
+    [m_RendererEncoder endEncoding];
 }
 
 void CommandBufferMTL::DrawIndexed(const DrawIndexedDesc& drawIndexedDesc) {
     id<MTLBuffer> indexBuffer = m_CurrentIndexCmd.m_Buffer->GetHandle();
-    [m_RendererEncoder drawIndexedPrimitives: m_CurrentPipeline->m_primitiveType
+    [m_RendererEncoder drawIndexedPrimitives: m_CurrentPipeline->GetPrimitiveType()
                               indexCount: drawIndexedDesc.indexNum
                                indexType: m_CurrentIndexCmd.m_Type
                              indexBuffer: indexBuffer
@@ -251,7 +308,7 @@
     NSCAssert(!countBuffer, @"count buffer not supported");
     
     [m_RendererEncoder
-         drawPrimitives: m_CurrentPipeline->m_primitiveType
+         drawPrimitives: m_CurrentPipeline->GetPrimitiveType()
          indirectBuffer:((BufferMTL&)buffer).GetHandle()
          indirectBufferOffset: offset];
 }
@@ -272,10 +329,16 @@
     
 }
 void CommandBufferMTL::DispatchIndirect(const Buffer& buffer, uint64_t offset) {}
-void CommandBufferMTL::BeginQuery(const QueryPool& queryPool, uint32_t offset) {}
+void CommandBufferMTL::BeginQuery(const QueryPool& queryPool, uint32_t offset) {
+    
+}
 void CommandBufferMTL::EndQuery(const QueryPool& queryPool, uint32_t offset) {}
-void CommandBufferMTL::BeginAnnotation(const char* name) {}
-void CommandBufferMTL::EndAnnotation() {}
+void CommandBufferMTL::BeginAnnotation(const char* name) {
+    
+}
+void CommandBufferMTL::EndAnnotation() {
+    
+}
 void CommandBufferMTL::ClearStorageBuffer(const ClearStorageBufferDesc& clearDesc) {
     
 }
@@ -286,6 +349,12 @@
 void CommandBufferMTL::CopyBuffer(Buffer& dstBuffer, uint64_t dstOffset, const Buffer& srcBuffer, uint64_t srcOffset, uint64_t size) {
     const BufferMTL& src = (const BufferMTL&)srcBuffer;
     const BufferMTL& dst = (const BufferMTL&)dstBuffer;
+    
+    if(!m_BlitEncoder) {
+        EndCurrentEncoders();
+        m_BlitEncoder = [m_Handle blitCommandEncoder];
+    }
+    
     [m_BlitEncoder
      copyFromBuffer:src.GetHandle()
      sourceOffset:srcOffset
@@ -297,6 +366,11 @@
 void CommandBufferMTL::CopyTexture(Texture& dstTexture, const TextureRegionDesc* dstRegionDesc, const Texture& srcTexture, const TextureRegionDesc* srcRegionDesc) {
     const TextureMTL& src = (const TextureMTL&)srcTexture;
     const TextureMTL& dst = (const TextureMTL&)dstTexture;
+    if(!m_BlitEncoder) {
+        EndCurrentEncoders();
+        m_BlitEncoder = [m_Handle blitCommandEncoder];
+    }
+    
     bool isWholeResource = !dstRegionDesc && !srcRegionDesc;
     if (isWholeResource) {
         [m_BlitEncoder
@@ -331,6 +405,11 @@
     const BufferMTL& src = (const BufferMTL&)srcBuffer;
     const TextureMTL& dst = (const TextureMTL&)dstTexture;
     
+    if(!m_BlitEncoder) {
+        EndCurrentEncoders();
+        m_BlitEncoder = [m_Handle blitCommandEncoder];
+    }
+    
     const MTLSize sourceSize = MTLSizeMake(
                                        (dstRegionDesc.width == WHOLE_SIZE) ? dst.GetSize(0, dstRegionDesc.mipOffset) : dstRegionDesc.width,
                                        (dstRegionDesc.height == WHOLE_SIZE) ? dst.GetSize(1, dstRegionDesc.mipOffset) : dstRegionDesc.height,
@@ -352,6 +431,10 @@
     const TextureMTL& src = (const TextureMTL&)srcTexture;
     const BufferMTL& dst = (const BufferMTL&)dstBuffer;
     
+    if(!m_BlitEncoder) {
+        EndCurrentEncoders();
+        m_BlitEncoder = [m_Handle blitCommandEncoder];
+    }
     
     const MTLSize sourceSize = MTLSizeMake(
                                        (srcRegionDesc.width == WHOLE_SIZE) ? src.GetSize(0, srcRegionDesc.mipOffset) : srcRegionDesc.width,
@@ -384,8 +467,14 @@
     
 }
 
-void Create(id<MTLCommandBuffer> cmd);
 
+void CommandBufferMTL::Create(const struct CommandQueueMTL* queue) {
+    m_CommandQueue = queue;
+    MTLCommandBufferDescriptor* pDesc = [[MTLCommandBufferDescriptor alloc] init];
+    pDesc.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
+    m_Handle = [m_CommandQueue->GetHandle() commandBufferWithDescriptor: pDesc];
+
+}
 
 
 #include "CommandBufferMTL.hpp"
diff --git a/Source/Metal/CommandQueueMTL.h b/Source/Metal/CommandQueueMTL.h
index 0486e73..8e0206d 100644
--- a/Source/Metal/CommandQueueMTL.h
+++ b/Source/Metal/CommandQueueMTL.h
@@ -7,12 +7,6 @@ namespace nri {
 
 struct DeviceMTL;
 
-NriBits(QueueBarrierBits, uint8_t,
-    NONE = 0,
-    BARRIER_FLAG_BUFFERS = NriBit(0),
-    BARRIER_FLAG_TEXTURES = NriBit(1),
-    BARRIER_FLAG_RENDERTARGETS = NriBit(2),
-    BARRIER_FLAG_FENCE = NriBit(3));
 
 struct CommandQueueMTL {
 
@@ -21,10 +15,6 @@ struct CommandQueueMTL {
     }
 
     ~CommandQueueMTL();
-    
-    inline operator id<MTLCommandQueue>() const {
-        return m_Handle;
-    }
 
     inline DeviceMTL& GetDevice() const {
         return m_Device;
@@ -38,14 +28,17 @@ struct CommandQueueMTL {
         return m_Lock;
     }
     
+    inline id<MTLCommandQueue> GetHandle() const {
+        return m_Handle;
+    }
+    
         
     void SetDebugName(const char* name);
     void Submit(const QueueSubmitDesc& queueSubmitDesc, const SwapChain* swapChain);
     Result WaitForIdle();
    
     Result Create(CommandQueueType type);
-    QueueBarrierBits m_BarrierBits = QueueBarrierBits::NONE;
-
+    
 private:
     DeviceMTL& m_Device;
     CommandQueueType m_Type = CommandQueueType(-1);
diff --git a/Source/Metal/CommandQueueMTL.mm b/Source/Metal/CommandQueueMTL.mm
index 6ad2415..47065bc 100644
--- a/Source/Metal/CommandQueueMTL.mm
+++ b/Source/Metal/CommandQueueMTL.mm
@@ -29,12 +29,9 @@
 void CommandQueueMTL::Submit(const QueueSubmitDesc& queueSubmitDesc, const SwapChain* swapChain) {
     
     for(uint32_t i = 0; i < queueSubmitDesc.commandBufferNum; i++) {
-        id<MTLCommandBuffer> cmd = *(struct CommandBufferMTL*)queueSubmitDesc.commandBuffers[i];
-        [cmd commit];
-        
+        const struct CommandBufferMTL& cmd = (const CommandBufferMTL&)queueSubmitDesc.commandBuffers[i];
+        [cmd.GetHandle() commit];
     }
-    
-    
 }
 
 
diff --git a/Source/Metal/ConversionMTL.h b/Source/Metal/ConversionMTL.h
index f8d1e8b..e321bcd 100644
--- a/Source/Metal/ConversionMTL.h
+++ b/Source/Metal/ConversionMTL.h
@@ -168,6 +168,8 @@ constexpr MTLDataType GetDescriptorType(DescriptorType type) {
     return DESCRIPTOR_TYPES[(size_t)type];
 }
 
+
+
 constexpr std::array<MTLCompareFunction, (size_t)CompareFunc::MAX_NUM> COMPARE_OP = {
     MTLCompareFunctionNever,            // NONE
     MTLCompareFunctionAlways,           // ALWAYS
diff --git a/Source/Metal/DescriptorMTL.h b/Source/Metal/DescriptorMTL.h
index b5a7447..c7594f3 100644
--- a/Source/Metal/DescriptorMTL.h
+++ b/Source/Metal/DescriptorMTL.h
@@ -19,7 +19,6 @@ enum class DescriptorTypeMTL {
   BUFFER_VIEW
 };
 
-
 struct DescriptorMTL {
 public:
     inline DescriptorMTL (DeviceMTL& device)
@@ -52,6 +51,8 @@ struct DescriptorMTL {
 
 private:
     
+    void EndCurrentEncoders(bool forceBarrier);
+    
     DeviceMTL& m_Device;
     DescriptorTypeMTL m_Type = DescriptorTypeMTL::NONE;
     id<MTLTexture> m_Texture;
diff --git a/Source/Metal/DeviceMTL.h b/Source/Metal/DeviceMTL.h
index eac2c87..2a5b307 100644
--- a/Source/Metal/DeviceMTL.h
+++ b/Source/Metal/DeviceMTL.h
@@ -48,6 +48,9 @@ struct DeviceMTL final : public DeviceBase {
     }
     
     
+    id<MTLRenderPipelineState> GetClearPipeline();
+    
+    
     void Destruct() override;
     Result FillFunctionTable(CoreInterface& table) const override;
     Result FillFunctionTable(HelperInterface& table) const override;
diff --git a/Source/Metal/DeviceMTL.mm b/Source/Metal/DeviceMTL.mm
index 244faea..fc359a7 100644
--- a/Source/Metal/DeviceMTL.mm
+++ b/Source/Metal/DeviceMTL.mm
@@ -82,6 +82,11 @@ static uint32_t GetEntryProperty(io_registry_entry_t entry, CFStringRef property
 //}
 
 
+id<MTLRenderPipelineState> DeviceMTL::GetClearPipeline() {
+    
+}
+
+
 void DeviceMTL::GetMemoryDesc(const BufferDesc& bufferDesc, MemoryLocation memoryLocation, MemoryDesc& memoryDesc) const {
     MemoryTypeInfo memoryTypeInfo;
     memoryTypeInfo.options = DEFAULT_MEMORY_RESOURCE_OPTION_MEMORY_LOCATION[(size_t)memoryLocation];
diff --git a/Source/Metal/ImplMTL.mm b/Source/Metal/ImplMTL.mm
index 0c8b311..384cb0c 100644
--- a/Source/Metal/ImplMTL.mm
+++ b/Source/Metal/ImplMTL.mm
@@ -68,6 +68,7 @@ static Result NRI_CALL CreateCommandBuffer(CommandAllocator& commandAllocator, C
 
 static void NRI_CALL ResetCommandAllocator(CommandAllocator& commandAllocator) {
    // ((CommandAllocatorVK&)commandAllocator).Reset();
+    
 }
 
 static void NRI_CALL SetCommandBufferDebugName(CommandBuffer& commandBuffer, const char* name) {
@@ -358,7 +359,7 @@ static Result NRI_CALL AllocateDescriptorSets(DescriptorPool& descriptorPool, co
     table.GetTextureMemoryDesc = ::GetTextureMemoryDesc;
     //table.GetCommandQueue = ::GetCommandQueue;
     //table.CreateCommandAllocator = ::CreateCommandAllocator;
-    //table.CreateCommandBuffer = ::CreateCommandBuffer;
+    table.CreateCommandBuffer = ::CreateCommandBuffer;
     //table.CreateDescriptorPool = ::CreateDescriptorPool;
     //table.CreateBuffer = ::CreateBuffer;
     table.CreateTexture = ::CreateTexture;
diff --git a/Source/Metal/PipelineLayoutMTL.h b/Source/Metal/PipelineLayoutMTL.h
index a783db6..e6416b5 100644
--- a/Source/Metal/PipelineLayoutMTL.h
+++ b/Source/Metal/PipelineLayoutMTL.h
@@ -19,6 +19,7 @@ struct PipelineLayoutMTL {
         , m_DescriptorSetRangeDescs(device.GetStdAllocator())
         , m_DynamicConstantBufferDescs(device.GetStdAllocator())
         , m_DescriptorSetLayouts(device.GetStdAllocator())
+        , m_PushBindingConstants(device.GetStdAllocator())
     {
     }
 
@@ -32,6 +33,10 @@ struct PipelineLayoutMTL {
         return &m_DescriptorSetLayouts[setIndex];
     }
 
+    const inline struct RootConstantDesc* GetPushBinding(uint32_t index) {
+        return &m_PushBindingConstants[index];
+    }
+    
     Result Create(const PipelineLayoutDesc& pipelineLayoutDesc);
    
 
@@ -42,6 +47,7 @@ struct PipelineLayoutMTL {
     Vector<DescriptorRangeDesc> m_DescriptorSetRangeDescs;
     Vector<DynamicConstantBufferDesc> m_DynamicConstantBufferDescs;
     Vector<DescriptorSetLayout> m_DescriptorSetLayouts;
+    Vector<RootConstantDesc> m_PushBindingConstants;
 
 };
 
diff --git a/Source/Metal/PipelineLayoutMTL.mm b/Source/Metal/PipelineLayoutMTL.mm
index ebb6287..fa88fad 100644
--- a/Source/Metal/PipelineLayoutMTL.mm
+++ b/Source/Metal/PipelineLayoutMTL.mm
@@ -22,6 +22,7 @@
     m_HasVariableDescriptorNum.reserve(pipelineLayoutDesc.descriptorSetNum);
     m_DescriptorSetRangeDescs.reserve(rangeNum);
     m_DynamicConstantBufferDescs.reserve(dynamicConstantBufferNum);
+    m_PushBindingConstants.resize(pipelineLayoutDesc.rootConstantNum);
     
     for (uint32_t i = 0; i < pipelineLayoutDesc.descriptorSetNum; i++) {
         const DescriptorSetDesc& descriptorSetDesc = pipelineLayoutDesc.descriptorSets[i];
@@ -71,6 +72,13 @@
         m_DescriptorSetLayouts[i].m_ArgumentDescriptors = argumentDescriptors;
 
     }
+    
+    
+    for (uint32_t i = 0; i < pipelineLayoutDesc.rootConstantNum; i++) {
+        const RootConstantDesc& pushConstantDesc = pipelineLayoutDesc.rootConstants[i];
+        m_PushBindingConstants[i] = pushConstantDesc;
+    }
+
     return Result::SUCCESS;
 }
 
diff --git a/Source/Metal/PipelineMTL.h b/Source/Metal/PipelineMTL.h
index 4351886..e6480fd 100644
--- a/Source/Metal/PipelineMTL.h
+++ b/Source/Metal/PipelineMTL.h
@@ -20,20 +20,34 @@ struct PipelineMTL {
     }
     ~PipelineMTL();
     
-    inline PipelineType GetPipelineType() {
-        return m_PipelineType;
-    }
-
     Result Create(const GraphicsPipelineDesc& graphicsPipelineDesc);
     Result Create(const ComputePipelineDesc& computePipelineDesc);
     Result Create(const RayTracingPipelineDesc& rayTracingPipelineDesc);
 
+    inline MTLPrimitiveType GetPrimitiveType() const {
+        return m_primitiveType;
+    }
+    
+    inline PipelineType GetPipelineType() const {
+        return m_PipelineType;
+    }
+
+    inline id<MTLComputePipelineState> GetComputePipeline() const {
+        return m_ComputePipeline;
+    }
+    
+    
+    inline id<MTLRenderPipelineState> GetGraphicsPipeline() const {
+        return m_GraphicsPipeline;
+    }
+    
+    
+private:
+    PipelineType m_PipelineType;
     MTLPrimitiveTopologyClass m_topologyClass;
     MTLPrimitiveType m_primitiveType;
     StageBits m_usedBits;
     
-private:
-    PipelineType m_PipelineType;
     union{
         id<MTLComputePipelineState> m_ComputePipeline = nil;
         id<MTLRenderPipelineState> m_GraphicsPipeline;