diff --git a/src/d3d12/d3d12_command_allocator.cpp b/src/d3d12/d3d12_command_allocator.cpp
index b60fd95ed..c7b47e1ee 100644
--- a/src/d3d12/d3d12_command_allocator.cpp
+++ b/src/d3d12/d3d12_command_allocator.cpp
@@ -75,6 +75,8 @@ MTLD3D12CommandAllocatorImpl::Initialize() {
   encoder_last = nullptr;
   encoder_count_ = 0;
 
+  icb_.clear();
+
   return S_OK;
 }
 
@@ -134,6 +136,188 @@ MTLD3D12CommandAllocatorImpl::Reset() {
   return Initialize();
 };
 
+IndirectComputeCommandData *
+MTLD3D12CommandAllocatorImpl::EncodeIndirectComputeCommand(MTLD3D12CommandSignature *pCmdSig, MTLD3D12ComputePipelineState *pPSO, size_t MaxCount) {
+  WMTIndirectCommandBufferInfo info;
+  info.inherit_buffers = !pCmdSig->UpdateRootArguments;
+  info.inherit_pso = 1;
+  info.inherit_cull_mode = 0;
+  info.inherit_fill_mode = 0;
+  info.inherit_front_facing = 0;
+  info.inherit_depth_bias = 0;
+  info.inherit_depth_clip_mode = 0;
+  info.inherit_depth_stencil_state = 0;
+  info.support_color_attachment_mapping = 0;
+  info.support_dynamic_attribute_stride = 0;
+  info.support_ray_tracing = 0;
+  info.type = WMTIndirectCommandTypeConcurrentDispatch;
+  info.max_vertex_buffer_binding = 0;
+  info.max_fragment_buffer_binding = 0;
+  info.max_object_buffer_binding = 0;
+  info.max_mesh_buffer_binding = 0;
+  info.max_kernel_buffer_binding = 31;
+  info.max_kernel_threadgroup_memory_binding = 0;
+  info.max_object_threadgroup_memory_binding = 0;
+  info.gpu_resource_id = 0;
+
+  auto icb = device_->GetMTLDevice().newIndirectCommandBuffer(info, MaxCount, WMTResourceStorageModeShared);
+
+  auto [Ptr, Offset] = AllocateGPUHeap(sizeof(IndirectComputeCommandData), 16);
+
+  auto data = reinterpret_cast<IndirectComputeCommandData *>(Ptr);
+
+  data->cmd_buf = info.gpu_resource_id;
+  data->max_count = MaxCount;
+  data->tgsize_x = pPSO->threadgroup_size.width;
+  data->tgsize_y = pPSO->threadgroup_size.height;
+  data->tgsize_z = pPSO->threadgroup_size.depth;
+
+  {
+    // populated outside
+    data->max_count_buffer = 0;
+    data->argument_buffer = 0;
+    data->rootsig_qwords = 0;
+    data->rootsig_qwords_stride = 0;
+    data->static_samplers = 0;
+  }
+
+  {
+    /**
+     * TODO: move these out?
+     */
+
+    auto &cmd_use_icb = EncodeComputeCommand<wmtcmd_compute_useresource>();
+    cmd_use_icb.type = WMTComputeCommandUseResource;
+    cmd_use_icb.usage = WMTResourceUsageRead | WMTResourceUsageWrite;
+    cmd_use_icb.resource = icb;
+
+    auto &cmd_setpso_res = EncodeComputeCommand<wmtcmd_compute_setpso>();
+    cmd_setpso_res.type = WMTComputeCommandSetPSO;
+    cmd_setpso_res.pso = pCmdSig->compute_resolver;
+    cmd_setpso_res.threadgroup_size = {1, 1, 1};
+
+    auto &cmd_argbuf_res = EncodeComputeCommand<wmtcmd_compute_setbuffer>();
+    cmd_argbuf_res.type = WMTComputeCommandSetBuffer;
+    cmd_argbuf_res.buffer = gpu_heap_buffer_;
+    cmd_argbuf_res.offset = Offset;
+    cmd_argbuf_res.index = 30;
+
+    auto &cmd_dispatch_res = EncodeComputeCommand<wmtcmd_compute_dispatch>();
+    cmd_dispatch_res.type = WMTComputeCommandDispatch;
+    cmd_dispatch_res.size = {1, 1, 1};
+
+    auto &cmd_setpso = EncodeComputeCommand<wmtcmd_compute_setpso>();
+    cmd_setpso.type = WMTComputeCommandSetPSO;
+    cmd_setpso.pso = pPSO->pso;
+    cmd_setpso.threadgroup_size = pPSO->threadgroup_size; // not really used
+  }
+
+  auto &cmd = EncodeComputeCommand<wmtcmd_compute_executecommands>();
+  cmd.type = WMTComputeCommandExecuteCommandsInBuffer;
+  cmd.indirect_command_buffer = icb;
+  cmd.location = 0;
+  cmd.length = MaxCount;
+
+  icb_.push_back(std::move(icb));
+
+  return data;
+}
+
+IndirectRenderCommandData *
+MTLD3D12CommandAllocatorImpl::EncodeIndirectRenderCommand(
+    MTLD3D12CommandSignature *pCmdSig, MTLD3D12GraphicsPipelineState *pPSO, size_t MaxCount
+) {
+  WMTIndirectCommandBufferInfo info;
+  info.inherit_buffers = !(pCmdSig->UpdateVertexBuffers || pCmdSig->UpdateIndexBuffer || pCmdSig->UpdateRootArguments);
+  info.inherit_pso = 1;
+  info.inherit_cull_mode = 1;
+  info.inherit_fill_mode = 1;
+  info.inherit_front_facing = 1;
+  info.inherit_depth_bias = 1;
+  info.inherit_depth_clip_mode = 1;
+  info.inherit_depth_stencil_state = 1;
+  info.support_color_attachment_mapping = 0;
+  info.support_dynamic_attribute_stride = 0;
+  info.support_ray_tracing = 0;
+  info.type = pCmdSig->CommandType == D3D12_INDIRECT_ARGUMENT_TYPE_DRAW_INDEXED ? WMTIndirectCommandTypeDrawIndexed
+                                                                                : WMTIndirectCommandTypeDraw;
+  info.max_vertex_buffer_binding = 31;
+  info.max_fragment_buffer_binding = 31;
+  info.max_object_buffer_binding = 0;
+  info.max_mesh_buffer_binding = 0;
+  info.max_kernel_buffer_binding = 0;
+  info.max_kernel_threadgroup_memory_binding = 0;
+  info.max_object_threadgroup_memory_binding = 0;
+  info.gpu_resource_id = 0;
+
+  auto icb = device_->GetMTLDevice().newIndirectCommandBuffer(info, MaxCount, WMTResourceStorageModePrivate);
+
+  auto [Ptr, Offset] = AllocateGPUHeap(sizeof(IndirectRenderCommandData), 16);
+
+  auto data = reinterpret_cast<IndirectRenderCommandData *>(Ptr);
+
+  data->cmd_buf = info.gpu_resource_id;
+  data->max_count = MaxCount;
+
+  {
+    // populated outside
+    data->max_count_buffer = 0;
+    data->argument_buffer = 0;
+    data->rootsig_qwords = 0;
+    data->rootsig_qwords_stride = 0;
+    data->static_samplers = 0;
+    data->vertex_buffer = 0;
+    data->vertex_argbuf_stride = 0;
+    data->primitive_type = 0;
+    data->index_buffer = 0;
+    data->index_buffer_format = {};
+  }
+
+  {
+    /**
+     * TODO: move these out?
+     */
+
+    auto &cmd_use_icb = EncodeRenderCommand<wmtcmd_render_useresource>();
+    cmd_use_icb.type = WMTRenderCommandUseResource;
+    cmd_use_icb.stages = WMTRenderStageVertex;
+    cmd_use_icb.usage = WMTResourceUsageRead | WMTResourceUsageWrite;
+    cmd_use_icb.resource = icb;
+
+    auto &cmd_setpso_res = EncodeRenderCommand<wmtcmd_render_setpso>();
+    cmd_setpso_res.type = WMTRenderCommandSetPSO;
+    cmd_setpso_res.pso = pCmdSig->render_resolver;
+
+    auto &cmd_argbuf_res = EncodeRenderCommand<wmtcmd_render_setbuffer>();
+    cmd_argbuf_res.type = WMTRenderCommandSetVertexBuffer;
+    cmd_argbuf_res.buffer = gpu_heap_buffer_;
+    cmd_argbuf_res.offset = Offset;
+    cmd_argbuf_res.index = 30;
+
+    auto &cmd_draw_res = EncodeRenderCommand<wmtcmd_render_draw>();
+    cmd_draw_res.type = WMTRenderCommandDraw;
+    cmd_draw_res.primitive_type = WMTPrimitiveTypePoint;
+    cmd_draw_res.vertex_start = 0;
+    cmd_draw_res.vertex_count = 1;
+    cmd_draw_res.base_instance = 0;
+    cmd_draw_res.instance_count = 1;
+
+    auto &cmd_setpso = EncodeRenderCommand<wmtcmd_render_setpso>();
+    cmd_setpso.type = WMTRenderCommandSetPSO;
+    cmd_setpso.pso = pPSO->pso;
+  }
+
+  auto &cmd = EncodeRenderCommand<wmtcmd_render_executecommands>();
+  cmd.type = WMTRenderCommandExecuteCommandsInBuffer;
+  cmd.indirect_command_buffer = icb;
+  cmd.location = 0;
+  cmd.length = MaxCount;
+
+  icb_.push_back(std::move(icb));
+
+  return data;
+}
+
 template <>
 WMT::Reference<WMT::ComputePipelineState>
 SimpleCommandContext<MTLD3D12CommandAllocatorImpl>::getComputePipeline(std::string name) {
diff --git a/src/d3d12/d3d12_command_allocator.hpp b/src/d3d12/d3d12_command_allocator.hpp
index 2096f7bb7..0b0a8d065 100644
--- a/src/d3d12/d3d12_command_allocator.hpp
+++ b/src/d3d12/d3d12_command_allocator.hpp
@@ -38,6 +38,34 @@ ptr_add(const void *const p, const std::uintptr_t &amount) noexcept {
   return reinterpret_cast<void *>(reinterpret_cast<std::uintptr_t>(p) + amount);
 }
 
+struct IndirectComputeCommandData {
+  uint64_t cmd_buf;
+  uint64_t max_count;
+  uint64_t max_count_buffer;
+  uint64_t argument_buffer;
+  uint64_t static_samplers;
+  uint64_t rootsig_qwords;
+  uint32_t rootsig_qwords_stride;
+  uint32_t tgsize_x;
+  uint32_t tgsize_y;
+  uint32_t tgsize_z;
+};
+
+struct IndirectRenderCommandData {
+  uint64_t cmd_buf;
+  uint64_t max_count;
+  uint64_t max_count_buffer;
+  uint64_t argument_buffer;
+  uint64_t static_samplers;
+  uint64_t rootsig_qwords;
+  uint32_t rootsig_qwords_stride;
+  uint32_t primitive_type;
+  uint64_t vertex_buffer;
+  uint64_t index_buffer;
+  DXGI_FORMAT index_buffer_format;
+  uint32_t vertex_argbuf_stride;
+};
+
 class MTLD3D12CommandAllocatorImpl : public MTLD3D12Pageable<MTLD3D12CommandAllocator> {
   friend class MTLD3D12GraphicsCommandListImpl;
   friend struct SimpleCommandContext<MTLD3D12CommandAllocatorImpl>;
@@ -58,6 +86,7 @@ class MTLD3D12CommandAllocatorImpl : public MTLD3D12Pageable<MTLD3D12CommandAllo
 
   small_vector<EncoderData, 64> encoder_lists_;
 
+  small_vector<WMT::Reference<WMT::IndirectCommandBuffer>, 4> icb_;
 
   ClearUAV<MTLD3D12CommandAllocatorImpl> clear_uav_;
 
@@ -186,6 +215,10 @@ class MTLD3D12CommandAllocatorImpl : public MTLD3D12Pageable<MTLD3D12CommandAllo
     assert(gpu_heap_offset_ < kGPUHeapSize);
     return {ptr_add(gpu_heap_, aligned), aligned};
   }
+
+  IndirectComputeCommandData *EncodeIndirectComputeCommand(MTLD3D12CommandSignature *pCmdSig, MTLD3D12ComputePipelineState *pPSO, size_t MaxCount);
+
+  IndirectRenderCommandData *EncodeIndirectRenderCommand(MTLD3D12CommandSignature *pCmdSig, MTLD3D12GraphicsPipelineState *pPSO, size_t MaxCount);
 };
 
 } // namespace dxmt
\ No newline at end of file
diff --git a/src/d3d12/d3d12_command_list.cpp b/src/d3d12/d3d12_command_list.cpp
index 30ee2378e..1d3642a9f 100644
--- a/src/d3d12/d3d12_command_list.cpp
+++ b/src/d3d12/d3d12_command_list.cpp
@@ -30,6 +30,8 @@ enum class DirtyState {
   ScissorRect,
   ComputeRootArguments,
   ComputeRootSignature,
+  BlendFactor,
+  StencilRef,
 };
 
 enum class DrawCallStatus {
@@ -145,6 +147,9 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
   Com<MTLD3D12RootSignature, false> rootsig_compute_;
   uint64_t rootarg_compute_staging_[64];
 
+  FLOAT blend_factor_[4];
+  UINT8 stencil_ref_;
+
 public:
   MTLD3D12GraphicsCommandListImpl(MTLD3D12Device *pDevice) : MTLD3D12DeviceChild<MTLD3D12GraphicsCommandList>(pDevice) {}
 
@@ -178,6 +183,12 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
     num_scissors = {};
     memset(scissors, 0, sizeof(scissors));
 
+    blend_factor_[0] = 1.0f;
+    blend_factor_[1] = 1.0f;
+    blend_factor_[2] = 1.0f;
+    blend_factor_[3] = 1.0f;
+    stencil_ref_ = 0;
+
     rootsig_graphics_ = nullptr;
     memset(rootarg_graphics_staging_, 0, sizeof(rootarg_graphics_staging_));
 
@@ -283,7 +294,7 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
   }
 
   DrawCallStatus
-  PreDraw() {
+  PreDraw(bool SkipResourceBinding = false) {
     if (!allocator_->encoder_current || allocator_->encoder_current->type != EncoderType::Render) {
 
       allocator_->InvalidateCurrentPass();
@@ -355,16 +366,16 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       }
       dirty_state_.set(DirtyState::VertexBuffer, DirtyState::GraphicsRootArguments, DirtyState::GraphicsRootSignature);
       dirty_state_.set(DirtyState::Viewport, DirtyState::ScissorRect);
+      dirty_state_.set(DirtyState::BlendFactor, DirtyState::StencilRef);
     }
     if (dirty_state_.test(DirtyState::VertexBuffer)) {
       EncodeVertexBuffers();
       dirty_state_.clr(DirtyState::VertexBuffer);
     }
 
-    if (dirty_state_.test(DirtyState::GraphicsRootArguments)) {
+    if (dirty_state_.test(DirtyState::GraphicsRootArguments) && !SkipResourceBinding) {
       if (rootsig_graphics_) {
-        auto [Ptr, Offset] = allocator_->AllocateGPUHeap(sizeof(uint64_t) * rootsig_graphics_->UploadQwords, 64);
-        memcpy(Ptr, rootarg_graphics_staging_, rootsig_graphics_->UploadQwords * sizeof(uint64_t));
+        auto Offset = EncodeRootArgument(rootsig_graphics_.ptr(), rootarg_graphics_staging_);
         auto &cmd_vsargbuf = allocator_->EncodeRenderCommand<wmtcmd_render_setbuffer>();
         cmd_vsargbuf.type = WMTRenderCommandSetVertexBuffer;
         cmd_vsargbuf.buffer = allocator_->gpu_heap_buffer_;
@@ -379,11 +390,9 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       dirty_state_.clr(DirtyState::GraphicsRootArguments);
     }
 
-    if (dirty_state_.test(DirtyState::GraphicsRootSignature)) {
+    if (dirty_state_.test(DirtyState::GraphicsRootSignature) && !SkipResourceBinding) {
       if (rootsig_graphics_) {
-        auto static_sampler_encode_size = sizeof(uint64_t) * rootsig_graphics_->NumStaticSamplers * 4;
-        auto [Ptr, Offset] = allocator_->AllocateGPUHeap(static_sampler_encode_size, 64);
-        memcpy(Ptr, rootsig_graphics_->EncodedStaticSamplers, static_sampler_encode_size);
+        auto Offset = EncodeStaticSamplers(rootsig_graphics_.ptr());
         auto &cmd_vsargbuf = allocator_->EncodeRenderCommand<wmtcmd_render_setbuffer>();
         cmd_vsargbuf.type = WMTRenderCommandSetVertexBuffer;
         cmd_vsargbuf.buffer = allocator_->gpu_heap_buffer_;
@@ -432,6 +441,24 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       cmd.rect_count = num_viewports;
       dirty_state_.clr(DirtyState::ScissorRect);
     }
+
+    if (dirty_state_.test(DirtyState::BlendFactor)) {
+      auto &cmd = allocator_->EncodeRenderCommand<wmtcmd_render_setblendcolor>();
+      cmd.type = WMTRenderCommandSetBlendFactor;
+      cmd.red = blend_factor_[0];
+      cmd.green = blend_factor_[1];
+      cmd.blue = blend_factor_[2];
+      cmd.alpha = blend_factor_[3];
+      dirty_state_.clr(DirtyState::BlendFactor);
+    }
+
+    if (dirty_state_.test(DirtyState::StencilRef)) {
+      auto &cmd = allocator_->EncodeRenderCommand<wmtcmd_render_setstencilref>();
+      cmd.type = WMTRenderCommandSetStencilRef;
+      cmd.stencil_ref = stencil_ref_;
+      dirty_state_.clr(DirtyState::StencilRef);
+    }
+
     return DrawCallStatus::Ordinary;
   }
 
@@ -478,8 +505,27 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
     cmd_draw.base_instance = StartInstanceLocation;
   };
 
+  uint64_t
+  EncodeRootArgument(MTLD3D12RootSignature *pRootSig, uint64_t const pStaging[64], UINT Count = 1) {
+    auto [Ptr, Offset] = allocator_->AllocateGPUHeap(sizeof(uint64_t) * pRootSig->UploadQwords * Count, 64);
+    for (unsigned i = 0; i < Count; i++)
+      memcpy(
+          reinterpret_cast<uint64_t *>(Ptr) + i * pRootSig->UploadQwords, pStaging,
+          pRootSig->UploadQwords * sizeof(uint64_t)
+      );
+    return Offset;
+  }
+
+  uint64_t
+  EncodeStaticSamplers(MTLD3D12RootSignature *pRootSig) {
+    auto static_sampler_encode_size = sizeof(uint64_t) * pRootSig->NumStaticSamplers * 4;
+    auto [Ptr, Offset] = allocator_->AllocateGPUHeap(static_sampler_encode_size, 64);
+    memcpy(Ptr, pRootSig->EncodedStaticSamplers, static_sampler_encode_size);
+    return Offset;
+  }
+
   bool
-  PreDispatch() {
+  PreDispatch(bool SkipResourceBinding = false) {
     if (!allocator_->encoder_current || allocator_->encoder_current->type != EncoderType::Compute) {
       allocator_->InvalidateCurrentPass();
       auto compute = allocator_->AllocatePass<ComputeEncoderData>();
@@ -496,10 +542,9 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       }
     }
 
-    if (dirty_state_.test(DirtyState::ComputeRootArguments)) {
+    if (dirty_state_.test(DirtyState::ComputeRootArguments) && !SkipResourceBinding) {
       if (rootsig_compute_) {
-        auto [Ptr, Offset] = allocator_->AllocateGPUHeap(sizeof(uint64_t) * rootsig_compute_->UploadQwords, 64);
-        memcpy(Ptr, rootarg_compute_staging_, rootsig_compute_->UploadQwords * sizeof(uint64_t));
+        auto Offset = EncodeRootArgument(rootsig_compute_.ptr(), rootarg_compute_staging_);
         auto &cmd_argbuf = allocator_->EncodeComputeCommand<wmtcmd_compute_setbuffer>();
         cmd_argbuf.type = WMTComputeCommandSetBuffer;
         cmd_argbuf.buffer = allocator_->gpu_heap_buffer_;
@@ -509,11 +554,9 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       dirty_state_.clr(DirtyState::ComputeRootArguments);
     }
 
-    if (dirty_state_.test(DirtyState::ComputeRootSignature)) {
+    if (dirty_state_.test(DirtyState::ComputeRootSignature) && !SkipResourceBinding) {
       if (rootsig_compute_) {
-        auto static_sampler_encode_size = sizeof(uint64_t) * rootsig_compute_->NumStaticSamplers * 4;
-        auto [Ptr, Offset] = allocator_->AllocateGPUHeap(static_sampler_encode_size, 64);
-        memcpy(Ptr, rootsig_compute_->EncodedStaticSamplers, static_sampler_encode_size);
+        auto Offset = EncodeStaticSamplers(rootsig_compute_.ptr());
         auto &cmd_argbuf = allocator_->EncodeComputeCommand<wmtcmd_compute_setbuffer>();
         cmd_argbuf.type = WMTComputeCommandSetBuffer;
         cmd_argbuf.buffer = allocator_->gpu_heap_buffer_;
@@ -649,11 +692,15 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
         if (!src)
           return;
 
-        auto &cmd_cp = allocator_->EncodeBlitCommand<wmtcmd_blit_copy_from_texture_to_buffer>();
-        cmd_cp.type = WMTBlitCommandCopyFromTextureToBuffer;
+        auto src_planar_count = getPlanarCount(src->pixelFormat());
+        auto src_subresource_index = pSrc->SubresourceIndex / src_planar_count;
+        auto src_subresource_planar = pSrc->SubresourceIndex % src_planar_count;
+
+        auto &cmd_cp = allocator_->EncodeBlitCommand<wmtcmd_blit_copy_from_texture_to_buffer_withblitoption>();
+        cmd_cp.type = WMTBlitCommandCopyFromTextureToBufferWithBlitOption;
         cmd_cp.src = src->current()->texture();
-        cmd_cp.level = pSrc->SubresourceIndex % src->miplevelCount();
-        cmd_cp.slice = pSrc->SubresourceIndex / src->miplevelCount();
+        cmd_cp.level = src_subresource_index % src->miplevelCount();
+        cmd_cp.slice = src_subresource_index / src->miplevelCount();
         cmd_cp.origin = {0, 0, 0};
         cmd_cp.size = {
             pDst->PlacedFootprint.Footprint.Width, pDst->PlacedFootprint.Footprint.Height,
@@ -665,6 +712,9 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
         cmd_cp.bytes_per_image = src->textureType() == WMTTextureType3D
                                      ? pDst->PlacedFootprint.Footprint.Height * pDst->PlacedFootprint.Footprint.RowPitch
                                      : 0;
+        cmd_cp.options = (src_planar_count > 1) ? (src_subresource_planar ? WMTBlitOptionStencilFromDepthStencil
+                                                                          : WMTBlitOptionDepthFromDepthStencil)
+                                                : WMTBlitOptionNone;
       } else {
         // so it is buffer to buffer copy?
         IMPLEMENT_ME
@@ -788,9 +838,24 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
     dirty_state_.set(DirtyState::ScissorRect);
   };
 
-  void STDMETHODCALLTYPE OMSetBlendFactor(const FLOAT BlendFactors[4]) { IMPLEMENT_ME };
+  void STDMETHODCALLTYPE
+  OMSetBlendFactor(const FLOAT BlendFactors[4]) {
+    if (BlendFactors) {
+      memcpy(blend_factor_, BlendFactors, std::size(blend_factor_) * sizeof(blend_factor_[0]));
+    } else {
+      blend_factor_[0] = 1.0f;
+      blend_factor_[1] = 1.0f;
+      blend_factor_[2] = 1.0f;
+      blend_factor_[3] = 1.0f;
+    }
+    dirty_state_.set(DirtyState::BlendFactor);
+  };
 
-  void STDMETHODCALLTYPE OMSetStencilRef(UINT StencilRef) { IMPLEMENT_ME };
+  void STDMETHODCALLTYPE
+  OMSetStencilRef(UINT StencilRef) {
+    stencil_ref_ = StencilRef;
+    dirty_state_.set(DirtyState::StencilRef);
+  };
 
   void
   UpdateGraphicsPSO(MTLD3D12GraphicsPipelineState *pso_graphics) {
@@ -1255,7 +1320,63 @@ class MTLD3D12GraphicsCommandListImpl : public MTLD3D12DeviceChild<MTLD3D12Graph
       ID3D12CommandSignature *pCommandSignature, UINT MaxCommandCount, ID3D12Resource *pArgBuffer,
       UINT64 ArgBufferOffset, ID3D12Resource *pCountBuffer, UINT64 CountBufferOffset
   ) {
-    IMPLEMENT_ME
+    auto sig = static_cast<MTLD3D12CommandSignature *>(pCommandSignature);
+    auto arg_buffer = static_cast<MTLD3D12Resource *>(pArgBuffer);
+    if (!arg_buffer || !arg_buffer->buffer)
+      return;
+    auto ArgBufferAddress = arg_buffer->buffer->current()->gpuAddress() + ArgBufferOffset;
+    uint64_t CountBufferAddress = 0;
+    if (auto count_buffer = static_cast<MTLD3D12Resource *>(pCountBuffer)) {
+      if (!count_buffer->buffer)
+        return;
+      CountBufferAddress = count_buffer->buffer->current()->gpuAddress() + CountBufferOffset;
+    }
+    if (sig->CommandType == D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH) {
+      if (!PreDispatch(sig->UpdateRootArguments))
+        return;
+
+      auto cmd = allocator_->EncodeIndirectComputeCommand(sig, pso_compute_.ptr(), MaxCommandCount);
+      cmd->max_count_buffer = CountBufferAddress;
+      cmd->argument_buffer = ArgBufferAddress;
+
+      if (sig->UpdateRootArguments) {
+        cmd->rootsig_qwords = EncodeRootArgument(rootsig_compute_.ptr(), rootarg_compute_staging_, MaxCommandCount);
+        cmd->rootsig_qwords += allocator_->gpu_heap_buffer_address_;
+        cmd->rootsig_qwords_stride = rootsig_compute_->UploadQwords;
+        cmd->static_samplers = EncodeStaticSamplers(rootsig_compute_.ptr());
+        cmd->static_samplers += allocator_->gpu_heap_buffer_address_;
+      }
+
+      return;
+    }
+    WMTPrimitiveType primitive_type;
+    uint32_t cp_count;
+    if (!to_metal_primitive_type(topology_, primitive_type, cp_count))
+      return;
+    bool encode_binding = sig->UpdateRootArguments || sig->UpdateIndexBuffer || sig->UpdateVertexBuffers;
+    DrawCallStatus status = PreDraw(encode_binding);
+    if (status == DrawCallStatus::Invalid)
+      return;
+    if (status != DrawCallStatus::Ordinary) {
+      IMPLEMENT_ME // TODO: (potential) emulated pipeline
+    }
+
+    auto cmd = allocator_->EncodeIndirectRenderCommand(sig, pso_graphics_.ptr(), MaxCommandCount);
+    cmd->max_count_buffer = CountBufferAddress;
+    cmd->argument_buffer = ArgBufferAddress;
+    cmd->primitive_type = primitive_type;
+    cmd->index_buffer = index_buffer_address;
+    cmd->index_buffer_format = index_type == WMTIndexTypeUInt32 ? DXGI_FORMAT_R32_UINT : DXGI_FORMAT_R16_UINT;
+    if (!encode_binding)
+      return;
+    cmd->rootsig_qwords = EncodeRootArgument(rootsig_graphics_.ptr(), rootarg_graphics_staging_, MaxCommandCount);
+    cmd->rootsig_qwords += allocator_->gpu_heap_buffer_address_;
+    cmd->rootsig_qwords_stride = rootsig_graphics_->UploadQwords;
+    cmd->static_samplers = EncodeStaticSamplers(rootsig_graphics_.ptr());
+    cmd->static_samplers += allocator_->gpu_heap_buffer_address_;
+    auto [VBOffset, VBStride] = PopulateVertexBufferTable(MaxCommandCount);
+    cmd->vertex_buffer = allocator_->gpu_heap_buffer_address_ + VBOffset;
+    cmd->vertex_argbuf_stride = VBStride;
   };
 };
 
diff --git a/src/d3d12/d3d12_command_signature.cpp b/src/d3d12/d3d12_command_signature.cpp
index 1b2c094c9..0a046e28b 100644
--- a/src/d3d12/d3d12_command_signature.cpp
+++ b/src/d3d12/d3d12_command_signature.cpp
@@ -22,6 +22,72 @@
 
 namespace dxmt {
 
+constexpr auto kSharedHeader = R"(
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct dxmt_compute_command_data {
+  command_buffer cmd_buf;
+  ulong max_count;
+  device uint * max_count_buffer;
+  device char * argument_buffer;
+  device ulong * static_samplers;
+  device ulong * rootsig_qwords;
+  uint rootsig_qwords_stride;
+  packed_uint3 tgsize;
+};
+
+struct d3d12_draw_arguments {
+  uint vertex_count_per_instance;
+  uint instance_count;
+  uint start_vertex_location;
+  uint start_instance_location;
+};
+
+struct d3d12_draw_indexed_arguments {
+  uint index_count_per_instance;
+  uint instance_count;
+  uint start_index_location;
+  int base_vertex_location;
+  uint start_instance_location;
+};
+
+struct d3d12_vertex_buffer_view {
+  device void * buffer;
+  uint size_in_bytes;
+  uint stride_in_bytes;
+};
+
+struct d3d12_index_buffer_view {
+  device void * buffer;
+  uint size_in_bytes;
+  uint format;
+};
+
+struct dxmt_vertex_buffer {
+  device void * buffer;
+  uint stride;
+  uint length;
+};
+
+struct dxmt_render_command_data {
+  command_buffer cmd_buf;
+  ulong max_count;
+  device uint * max_count_buffer;
+  device char * argument_buffer;
+  device ulong * static_samplers;
+  device ulong * rootsig_qwords;
+  uint rootsig_qwords_stride;
+  uint primitive_type;
+  device char * vertex_buffer;
+  device void * index_buffer;
+  uint index_buffer_format;
+  uint vertex_argbuf_stride;
+};
+
+)";
+
 class MTLD3D12CommandSignatureImpl : public MTLD3D12Pageable<MTLD3D12CommandSignature> {
 
 public:
@@ -29,6 +95,260 @@ class MTLD3D12CommandSignatureImpl : public MTLD3D12Pageable<MTLD3D12CommandSign
 
   HRESULT
   Initialize(const D3D12_COMMAND_SIGNATURE_DESC *pDesc, ID3D12RootSignature *pRootSignature) {
+    std::stringstream source;
+    D3D12_INDIRECT_ARGUMENT_TYPE side_effect = ~(D3D12_INDIRECT_ARGUMENT_TYPE){};
+    UpdateRootArguments = false;
+    UpdateVertexBuffers = false;
+    uint32_t ib_index = ~-0u;
+
+    source << kSharedHeader;
+
+    source << "struct __attribute__ ((packed)) d3d12_arguments {\n";
+
+    for (unsigned i = 0; i < pDesc->NumArgumentDescs; i++) {
+      if (~side_effect != 0u)
+        return E_INVALIDARG;
+      auto &arg = pDesc->pArgumentDescs[i];
+      switch (arg.Type) {
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH: {
+        side_effect = arg.Type;
+        source << "packed_uint3 dispatch;\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DRAW: {
+        side_effect = arg.Type;
+        source << "d3d12_draw_arguments draw;\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DRAW_INDEXED: {
+        side_effect = arg.Type;
+        source << "d3d12_draw_indexed_arguments draw_indexed;\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_VERTEX_BUFFER_VIEW: {
+        UpdateVertexBuffers = true;
+        source << "d3d12_vertex_buffer_view vb_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_INDEX_BUFFER_VIEW: {
+        if (ib_index != ~0u)
+          return E_INVALIDARG;
+        ib_index = i;
+        source << "d3d12_index_buffer_view ib;\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT: {
+        UpdateRootArguments = true;
+        for (unsigned j = 0; j < arg.Constant.Num32BitValuesToSet; j++) {
+          source << "uint constant_" << i << "_" << j << ";\n";
+        }
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT_BUFFER_VIEW: {
+        UpdateRootArguments = true;
+        source << "ulong cb_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_SHADER_RESOURCE_VIEW: {
+        UpdateRootArguments = true;
+        source << "ulong srv_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_UNORDERED_ACCESS_VIEW: {
+        UpdateRootArguments = true;
+        source << "ulong uav_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH_RAYS:
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH_MESH:
+        ERR("D3D12CommandSignatuer: unsupported rays/mesh dispatch");
+        return E_NOTIMPL;
+      default:
+        return E_INVALIDARG;
+      }
+    }
+    source << "};\n\n";
+
+    if (~side_effect == 0)
+      return E_INVALIDARG;
+    bool is_compute = side_effect == D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH;
+    CommandType = side_effect;
+    UpdateIndexBuffer = ib_index != ~0u;
+
+    if (is_compute)
+      source
+          << "[[kernel]] void resolve_indirect_commands([[thread_position_in_grid]] uint x, constant dxmt_compute_command_data";
+    else
+      source << "[[vertex]] void resolve_indirect_commands(constant dxmt_render_command_data";
+    source << " &command_data [[buffer(30)]]) {\n";
+
+    if (is_compute)
+      source << "if (x !=0 ) return;\n";
+
+    source << "uint count = command_data.max_count_buffer ? "
+              "command_data.max_count_buffer[0] : command_data.max_count;\n";
+    source << "for (uint i = 0; i < command_data.max_count; i++) {\n";
+    source << "device d3d12_arguments& arg = reinterpret_cast<device d3d12_arguments *>("
+              "command_data.argument_buffer + i * "
+           << pDesc->ByteStride << ")[0];\n";
+    if (is_compute) {
+      source << "compute_command cmd(command_data.cmd_buf, i);\n";
+    } else {
+      source << "render_command cmd(command_data.cmd_buf, i);\n";
+    }
+    source << "cmd.reset();\n";
+    source << "if (i >= count) continue;\n";
+    source << "device ulong * rootsig_qwords = command_data.rootsig_qwords + "
+              "(i * command_data.rootsig_qwords_stride);\n";
+    if (!is_compute)
+      source << "device dxmt_vertex_buffer * vertex_buffer = "
+                "reinterpret_cast<device dxmt_vertex_buffer *>(command_data.vertex_buffer + "
+                "(i * command_data.vertex_argbuf_stride));\n";
+
+    if (UpdateRootArguments || UpdateVertexBuffers || UpdateIndexBuffer) {
+      if (!is_compute) {
+        source << "cmd.set_vertex_buffer(vertex_buffer," << SM50_BINDING_INDEX_VERTEX_BUFFER << ");\n";
+        source << "cmd.set_vertex_buffer(rootsig_qwords," << SM50_BINDING_INDEX_ROOT_ARGUMENTS << ");\n";
+        source << "cmd.set_vertex_buffer(command_data.static_samplers," << SM50_BINDING_INDEX_STATIC_SAMPLERS << ");\n";
+        source << "cmd.set_fragment_buffer(rootsig_qwords," << SM50_BINDING_INDEX_ROOT_ARGUMENTS << ");\n";
+        source << "cmd.set_fragment_buffer(command_data.static_samplers," << SM50_BINDING_INDEX_STATIC_SAMPLERS
+               << ");\n";
+      } else {
+        source << "cmd.set_kernel_buffer(rootsig_qwords, " << SM50_BINDING_INDEX_ROOT_ARGUMENTS << ");\n";
+        source << "cmd.set_kernel_buffer(command_data.static_samplers," << SM50_BINDING_INDEX_STATIC_SAMPLERS << ");\n";
+      }
+    }
+
+    for (unsigned i = 0; i < pDesc->NumArgumentDescs; i++) {
+      auto &arg = pDesc->pArgumentDescs[i];
+      switch (arg.Type) {
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DRAW: {
+        source << "cmd.draw_primitives((primitive_type)command_data.primitive_type, "
+                  "arg.draw.start_vertex_location, "
+                  "arg.draw.vertex_count_per_instance, arg.draw.instance_count, "
+                  "arg.draw.start_instance_location);\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DRAW_INDEXED: {
+        if (ib_index == ~0u) {
+          source << "bool ib32bit = command_data.index_buffer_format == 42;\n";
+          source << "device void* ib = command_data.index_buffer;\n";
+        }
+        source << "if (ib32bit) {\n";
+        source << "cmd.draw_indexed_primitives((primitive_type)command_data.primitive_type, "
+                  "arg.draw_indexed.index_count_per_instance, "
+                  "reinterpret_cast<device uint *>(ib) + arg.draw_indexed.start_index_location, "
+                  "arg.draw_indexed.instance_count, arg.draw_indexed.base_vertex_location, "
+                  "arg.draw_indexed.start_instance_location);\n";
+        source << "} else {\n";
+        source << "cmd.draw_indexed_primitives((primitive_type)command_data.primitive_type, "
+                  "arg.draw_indexed.index_count_per_instance, "
+                  "reinterpret_cast<device ushort *>(ib) + arg.draw_indexed.start_index_location, "
+                  "arg.draw_indexed.instance_count, arg.draw_indexed.base_vertex_location, "
+                  "arg.draw_indexed.start_instance_location);\n";
+        source << "}\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_DISPATCH: {
+        source << "cmd.concurrent_dispatch_threadgroups(arg.dispatch, command_data.tgsize);\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_VERTEX_BUFFER_VIEW: {
+        auto slot = arg.VertexBuffer.Slot;
+        source << "vertex_buffer[" << slot << "] = {arg.vb_" << i << ".buffer,arg.vb_" << i
+               << ".stride_in_bytes,arg.vb_" << i << ".size_in_bytes};\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_INDEX_BUFFER_VIEW: {
+        source << "bool ib32bit = arg.ib.format == 42;\n";
+        source << "device void* ib = arg.ib.buffer;\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT: {
+        auto parameter_index = arg.Constant.RootParameterIndex;
+        if (!pRootSignature)
+          return E_INVALIDARG;
+        auto rootsig = static_cast<MTLD3D12RootSignature *>(pRootSignature);
+        if (parameter_index >= rootsig->ParameterSlots)
+          return E_INVALIDARG;
+        auto offset = rootsig->SlotQwordOffsets[parameter_index];
+        for (unsigned j = 0; j < arg.Constant.Num32BitValuesToSet; j++) {
+          source << "reinterpret_cast<device uint *>(rootsig_qwords + " << offset << ")["
+                 << (j + arg.Constant.DestOffsetIn32BitValues) << "] = arg.constant_" << i << "_" << j << ";\n";
+        }
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT_BUFFER_VIEW: {
+        auto parameter_index = arg.ConstantBufferView.RootParameterIndex;
+        if (!pRootSignature)
+          return E_INVALIDARG;
+        auto rootsig = static_cast<MTLD3D12RootSignature *>(pRootSignature);
+        if (parameter_index >= rootsig->ParameterSlots)
+          return E_INVALIDARG;
+        auto offset = rootsig->SlotQwordOffsets[parameter_index];
+        source << "rootsig_qwords[" << offset << "] = arg.cb_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_SHADER_RESOURCE_VIEW: {
+        auto parameter_index = arg.ShaderResourceView.RootParameterIndex;
+        if (!pRootSignature)
+          return E_INVALIDARG;
+        auto rootsig = static_cast<MTLD3D12RootSignature *>(pRootSignature);
+        if (parameter_index >= rootsig->ParameterSlots)
+          return E_INVALIDARG;
+        auto offset = rootsig->SlotQwordOffsets[parameter_index];
+        source << "rootsig_qwords[" << offset << "] = arg.srv_" << i << ";\n";
+        break;
+      }
+      case D3D12_INDIRECT_ARGUMENT_TYPE_UNORDERED_ACCESS_VIEW: {
+        auto parameter_index = arg.UnorderedAccessView.RootParameterIndex;
+        if (!pRootSignature)
+          return E_INVALIDARG;
+        auto rootsig = static_cast<MTLD3D12RootSignature *>(pRootSignature);
+        if (parameter_index >= rootsig->ParameterSlots)
+          return E_INVALIDARG;
+        auto offset = rootsig->SlotQwordOffsets[parameter_index];
+        source << "rootsig_qwords[" << offset << "] = arg.uav_" << i << ";\n";
+        break;
+      }
+      default:
+        return E_INVALIDARG;
+      }
+    }
+
+    source << "}\n"
+              "};\n";
+
+    WMT::Reference<WMT::Error> err;
+    auto lib = device_->GetMTLDevice().newLibraryWithSource(source.view(), err);
+
+    if (!lib) {
+      ERR("Failed to compile command signature resolve shader: ", err.description().getUTF8String());
+      return E_FAIL;
+    }
+
+    auto function = lib.newFunction("resolve_indirect_commands");
+
+    if (!function) {
+      ERR("Failed to create command signature resolve shader");
+      return E_FAIL;
+    }
+
+    if (is_compute) {
+      compute_resolver = device_->GetMTLDevice().newComputePipelineState(function, err);
+    } else {
+      WMTRenderPipelineInfo info;
+      WMT::InitializeRenderPipelineInfo(info);
+      info.rasterization_enabled = false;
+      info.vertex_function = function;
+      render_resolver = device_->GetMTLDevice().newRenderPipelineState(info, err);
+    }
+
+    if (err) {
+      ERR("Failed to compile command signature resolve pso: ", err.description().getUTF8String());
+      return E_FAIL;
+    }
+
     return S_OK;
   };
 
diff --git a/src/d3d12/d3d12_descriptor_heap.cpp b/src/d3d12/d3d12_descriptor_heap.cpp
index 57be437ef..18e276040 100644
--- a/src/d3d12/d3d12_descriptor_heap.cpp
+++ b/src/d3d12/d3d12_descriptor_heap.cpp
@@ -59,6 +59,7 @@ struct ShaderVisibleDescriptorGPUStorage {
     UAVBufferGPUStorage UAVBuffer;
     SRVTexelBufferGPUStorage SRVTexelBuffer;
     SRVBufferGPUStorage SRVBuffer;
+    std::array<uint64_t, 4> ZeroFilled;
   };
 
   ShaderVisibleDescriptorGPUStorage();
@@ -302,6 +303,42 @@ class MTLD3D12DescriptorHeapImpl : public MTLD3D12Pageable<MTLD3D12DescriptorHea
     return S_OK;
   }
 
+  virtual HRESULT
+  AddShaderResourceView(UINT Index, D3D12_SHADER_RESOURCE_VIEW_DESC const *pDesc) {
+    if (Index >= descriptors_.size())
+      return E_INVALIDARG;
+    if (!pDesc)
+      return E_INVALIDARG;
+    /**
+     * TODO: support null descriptor properly (respect different view dimensions)
+     */
+    auto &cpu_storage = descriptors_[Index];
+    cpu_storage.type = ShaderVisibleDescriptorType::Null;
+    if (mapped_argument_buffer_) {
+      auto &gpu_storage = mapped_argument_buffer_[Index];
+      gpu_storage.ZeroFilled = {{}};
+    }
+    return S_OK;
+  }
+
+  virtual HRESULT
+  AddUnorderedAccessView(UINT Index, D3D12_UNORDERED_ACCESS_VIEW_DESC const *pDesc) {
+    if (Index >= descriptors_.size())
+      return E_INVALIDARG;
+    if (!pDesc)
+      return E_INVALIDARG;
+    /**
+     * TODO: support null descriptor properly (respect different view dimensions)
+     */
+    auto &cpu_storage = descriptors_[Index];
+    cpu_storage.type = ShaderVisibleDescriptorType::Null;
+    if (mapped_argument_buffer_) {
+      auto &gpu_storage = mapped_argument_buffer_[Index];
+      gpu_storage.ZeroFilled = {{}};
+    }
+    return S_OK;
+  }
+
   virtual ShaderVisibleDescriptorCPUStorage const &
   GetDescriptor(UINT Index) {
     return descriptors_[Index];
diff --git a/src/d3d12/d3d12_descriptor_heap.hpp b/src/d3d12/d3d12_descriptor_heap.hpp
index 3db3bf446..3d8a1f2e1 100644
--- a/src/d3d12/d3d12_descriptor_heap.hpp
+++ b/src/d3d12/d3d12_descriptor_heap.hpp
@@ -147,6 +147,10 @@ class MTLD3D12DescriptorHeap : public ID3D12DescriptorHeap {
 
   virtual HRESULT AddShaderResourceView(UINT Index, Buffer *Buffer, BufferSlice Slice) = 0;
 
+  virtual HRESULT AddShaderResourceView(UINT Index, D3D12_SHADER_RESOURCE_VIEW_DESC const *pDesc) = 0;
+
+  virtual HRESULT AddUnorderedAccessView(UINT Index, D3D12_UNORDERED_ACCESS_VIEW_DESC const *pDesc) = 0;
+
   virtual ShaderVisibleDescriptorCPUStorage const &GetDescriptor(UINT Index) = 0;
 
   virtual void CopyDescriptors(UINT From, MTLD3D12DescriptorHeap *pHeapTo, UINT DescriptorTo, UINT CopyCount) = 0;
diff --git a/src/d3d12/d3d12_device.cpp b/src/d3d12/d3d12_device.cpp
index 39413d0e6..2f630f067 100644
--- a/src/d3d12/d3d12_device.cpp
+++ b/src/d3d12/d3d12_device.cpp
@@ -332,8 +332,9 @@ class MTLD3D12DeviceImpl : public MTLD3D12Object<ComObject<MTLD3D12Device>> {
       ID3D12Resource *pResource, const D3D12_SHADER_RESOURCE_VIEW_DESC *pDesc, D3D12_CPU_DESCRIPTOR_HANDLE Descriptor
   ) {
     if (!pResource) {
-      // null descriptor
-      IMPLEMENT_ME
+      auto [Heap, Index] = GetShaderVisibleDescriptorHeap(this, Descriptor);
+      Heap->AddShaderResourceView(Index, pDesc);
+      return;
     }
     auto d3d12res = static_cast<MTLD3D12Resource *>(pResource);
     d3d12res->CreateShaderResourceView(pDesc, Descriptor);
@@ -345,8 +346,9 @@ class MTLD3D12DeviceImpl : public MTLD3D12Object<ComObject<MTLD3D12Device>> {
       D3D12_CPU_DESCRIPTOR_HANDLE Descriptor
   ) {
     if (!pResource) {
-      // null descriptor
-      IMPLEMENT_ME
+      auto [Heap, Index] = GetShaderVisibleDescriptorHeap(this, Descriptor);
+      Heap->AddUnorderedAccessView(Index, pDesc);
+      return;
     }
     auto d3d12res = static_cast<MTLD3D12Resource *>(pResource);
     d3d12res->CreateUnorderedAccessView(pCounter, pDesc, Descriptor);
diff --git a/src/d3d12/d3d12_device.hpp b/src/d3d12/d3d12_device.hpp
index 9d08a79df..007122f1a 100644
--- a/src/d3d12/d3d12_device.hpp
+++ b/src/d3d12/d3d12_device.hpp
@@ -108,6 +108,16 @@ class MTLD3D12RootSignature : public ID3D12RootSignature {
 
 class MTLD3D12CommandSignature : public ID3D12CommandSignature {
 public:
+  D3D12_INDIRECT_ARGUMENT_TYPE CommandType;
+  UINT UpdateRootArguments : 1;
+  UINT UpdateVertexBuffers : 1;
+  UINT UpdateIndexBuffer   : 1;
+
+  WMT::Reference<WMT::RenderPipelineState> render_resolver;
+  WMT::Reference<WMT::ComputePipelineState> compute_resolver;
+
+  virtual void AddRefPrivate() = 0;
+  virtual void ReleasePrivate() = 0;
 };
 
 class MTLD3D12QueryHeap : public ID3D12QueryHeap {
diff --git a/src/d3d12/d3d12_pipeline_compute.cpp b/src/d3d12/d3d12_pipeline_compute.cpp
index bd0c8e3f6..024629062 100644
--- a/src/d3d12/d3d12_pipeline_compute.cpp
+++ b/src/d3d12/d3d12_pipeline_compute.cpp
@@ -89,6 +89,7 @@ class MTLD3D12ComputePipelineStateImpl : public MTLD3D12Pageable<MTLD3D12Compute
       WMTComputePipelineInfo info;
       WMT::InitializeComputePipelineInfo(info);
       info.compute_function = cs_func;
+      info.support_indirect_command_buffers = true;
 
       pso = metal.newComputePipelineState(info, err);
       if (!pso) {
diff --git a/src/d3d12/d3d12_pipeline_graphics.cpp b/src/d3d12/d3d12_pipeline_graphics.cpp
index 269aa759b..97552d8d5 100644
--- a/src/d3d12/d3d12_pipeline_graphics.cpp
+++ b/src/d3d12/d3d12_pipeline_graphics.cpp
@@ -462,6 +462,7 @@ class MTLD3D12GraphicsPipelineStateImpl : public MTLD3D12Pageable<MTLD3D12Graphi
       }
 
       info.raster_sample_count = pDesc->SampleDesc.Count;
+      info.support_indirect_command_buffers = true;
 
       pso = metal.newRenderPipelineState(info, err);
 
diff --git a/src/winemetal/Metal.hpp b/src/winemetal/Metal.hpp
index 70e2960e3..1f51b6d48 100644
--- a/src/winemetal/Metal.hpp
+++ b/src/winemetal/Metal.hpp
@@ -294,6 +294,19 @@ class Buffer : public Resource {
   }
 };
 
+class Heap : public Allocation {
+public:
+  Reference<Buffer>
+  newBuffer(WMTBufferInfo &info, uint64_t offset = ~0ull) {
+    return Reference<Buffer>(MTLHeap_newBuffer(handle, &info, offset));
+  }
+
+  Reference<Texture>
+  newTexture(WMTTextureInfo &info, uint64_t offset = ~0ull) {
+    return Reference<Texture>(MTLHeap_newTexture(handle, &info, offset));
+  }
+};
+
 class SamplerState : public Object {
 public:
 };
@@ -740,6 +753,10 @@ class BinaryArchive : public Object {
   }
 };
 
+class IndirectCommandBuffer : public Resource {
+public:
+};
+
 class Device : public Object {
 public:
   uint64_t
@@ -818,6 +835,11 @@ class Device : public Object {
     return Reference<Library>(MTLDevice_newLibrary(handle, data, &error.handle));
   }
 
+  Reference<Library>
+  newLibraryWithSource(std::string_view view, Error &error) {
+    return Reference<Library>(MTLDevice_newLibraryWithSource(handle, view.data(), view.length(), &error.handle));
+  }
+
   Reference<ComputePipelineState>
   newComputePipelineState(const Function &compute_function, Error &error) {
     WMTComputePipelineInfo info;
@@ -828,6 +850,7 @@ class Device : public Object {
     info.binary_archives_for_lookup.set(nullptr);
     info.num_binary_archives_for_lookup = 0;
     info.fail_on_binary_archive_miss = false;
+    info.support_indirect_command_buffers = false;
     return Reference<ComputePipelineState>(MTLDevice_newComputePipelineState(handle, &info, &error.handle));
   }
 
@@ -841,6 +864,7 @@ class Device : public Object {
     info.binary_archives_for_lookup.set(nullptr);
     info.num_binary_archives_for_lookup = 0;
     info.fail_on_binary_archive_miss = false;
+    info.support_indirect_command_buffers = false;
     return Reference<ComputePipelineState>(MTLDevice_newComputePipelineState(handle, &info, &error.handle));
   }
 
@@ -879,6 +903,21 @@ class Device : public Object {
     return Reference<ResidencySet>(MTLDevice_newResidencySet(handle, init_capacity, &error.handle));
   }
 
+  Reference<Heap>
+  newHeap(const WMTHeapInfo &info) {
+    return Reference<Heap>(MTLDevice_newHeap(handle, &info));
+  }
+
+  WMTSizeAndAlign
+  heapBufferSizeAndAlign(uint64_t length, WMTResourceOptions options) {
+    return MTLDevice_heapBufferSizeAndAlign(handle, length, options);
+  }
+
+  WMTSizeAndAlign
+  heapTextureSizeAndAlign(WMTTextureInfo &info) {
+    return MTLDevice_heapTextureSizeAndAlign(handle, &info);
+  }
+
   uint64_t
   minimumLinearTextureAlignmentForPixelFormat(WMTPixelFormat format) {
     return MTLDevice_minimumLinearTextureAlignmentForPixelFormat(handle, format);
@@ -943,6 +982,11 @@ class Device : public Object {
   newCounterSampleBuffer(uint32_t sample_count, bool shared = true) {
     return Reference<CounterSampleBuffer>(MTLCounterSampleBuffer_newTimestampBuffer(handle, sample_count, shared));
   }
+
+  Reference<IndirectCommandBuffer>
+  newIndirectCommandBuffer(WMTIndirectCommandBufferInfo &info, uint64_t max_count, WMTResourceOptions options) {
+    return Reference<IndirectCommandBuffer>(MTLDevice_newIndirectCommandBuffer(handle, &info, max_count, options));
+  }
 };
 
 inline Reference<Array<Device>>
@@ -1106,6 +1150,7 @@ InitializeRenderPipelineInfo(WMTRenderPipelineInfo &info) {
   info.binary_archives_for_lookup.set(nullptr);
   info.num_binary_archives_for_lookup = 0;
   info.fail_on_binary_archive_miss = false;
+  info.support_indirect_command_buffers = false;
 }
 
 inline void
@@ -1117,6 +1162,7 @@ InitializeComputePipelineInfo(WMTComputePipelineInfo &info) {
   info.fail_on_binary_archive_miss = false;
   info.tgsize_is_multiple_of_sgwidth = false;
   info.immutable_buffers = 0;
+  info.support_indirect_command_buffers = false;
 }
 
 inline void
@@ -1153,6 +1199,7 @@ InitializeMeshRenderPipelineInfo(WMTMeshRenderPipelineInfo &info) {
   info.binary_archives_for_lookup.set(nullptr);
   info.num_binary_archives_for_lookup = 0;
   info.fail_on_binary_archive_miss = false;
+  info.support_indirect_command_buffers = false;
 }
 
 inline void
diff --git a/src/winemetal/unix/winemetal_unix.c b/src/winemetal/unix/winemetal_unix.c
index 0bfd7e289..cd7df600a 100644
--- a/src/winemetal/unix/winemetal_unix.c
+++ b/src/winemetal/unix/winemetal_unix.c
@@ -377,6 +377,7 @@ _MTLDevice_newComputePipelineState(void *obj) {
   NSError *err = NULL;
   descriptor.computeFunction = (id<MTLFunction>)info->compute_function;
   descriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth = info->tgsize_is_multiple_of_sgwidth;
+  descriptor.supportIndirectCommandBuffers = info->support_indirect_command_buffers;
   for (unsigned i = 0; i < 31; i++) {
     if (info->immutable_buffers & (1 << i))
       descriptor.buffers[i].mutability = MTLMutabilityImmutable;
@@ -561,6 +562,7 @@ _MTLDevice_newRenderPipelineState(void *obj) {
 
   descriptor.vertexFunction = (id<MTLFunction>)info->vertex_function;
   descriptor.fragmentFunction = (id<MTLFunction>)info->fragment_function;
+  descriptor.supportIndirectCommandBuffers = info->support_indirect_command_buffers;
 
   if (info->num_binary_archives_for_lookup && info->binary_archives_for_lookup.ptr)
     descriptor.binaryArchives = [NSArray arrayWithObjects:(id<MTLBinaryArchive> *)info->binary_archives_for_lookup.ptr
@@ -627,6 +629,7 @@ _MTLDevice_newMeshRenderPipelineState(void *obj) {
 
   descriptor.meshThreadgroupSizeIsMultipleOfThreadExecutionWidth = info->mesh_tgsize_is_multiple_of_sgwidth;
   descriptor.objectThreadgroupSizeIsMultipleOfThreadExecutionWidth = info->object_tgsize_is_multiple_of_sgwidth;
+  descriptor.supportIndirectCommandBuffers = info->support_indirect_command_buffers;
 
   MTLPipelineOption options = MTLPipelineOptionNone;
 #if __MAC_OS_X_VERSION_MAX_ALLOWED >= 150000
@@ -716,6 +719,21 @@ _MTLBlitCommandEncoder_encodeCommands(void *obj) {
           destinationBytesPerImage:body->bytes_per_image];
       break;
     }
+    case WMTBlitCommandCopyFromTextureToBufferWithBlitOption: {
+      struct wmtcmd_blit_copy_from_texture_to_buffer_withblitoption *body =
+          (struct wmtcmd_blit_copy_from_texture_to_buffer_withblitoption *)next;
+      [encoder copyFromTexture:(id<MTLTexture>)body->src
+                       sourceSlice:body->slice
+                       sourceLevel:body->level
+                      sourceOrigin:MTLOriginMake(body->origin.x, body->origin.y, body->origin.z)
+                        sourceSize:MTLSizeMake(body->size.width, body->size.height, body->size.depth)
+                          toBuffer:(id<MTLBuffer>)body->dst
+                 destinationOffset:body->offset
+            destinationBytesPerRow:body->bytes_per_row
+          destinationBytesPerImage:body->bytes_per_image
+                           options:(MTLBlitOption)body->options];
+      break;
+    }
     case WMTBlitCommandCopyFromTextureToTexture: {
       struct wmtcmd_blit_copy_from_texture_to_texture *body = (struct wmtcmd_blit_copy_from_texture_to_texture *)next;
       [encoder copyFromTexture:(id<MTLTexture>)body->src
@@ -757,6 +775,12 @@ _MTLBlitCommandEncoder_encodeCommands(void *obj) {
              destinationOffset:body->dst_offset];
       break;
     }
+    case WMTBlitCommandResetCommandsInBuffer: {
+      struct wmtcmd_blit_resetcommands *body = (struct wmtcmd_blit_resetcommands *)next;
+      [encoder resetCommandsInBuffer:(id<MTLIndirectCommandBuffer>)body->indirect_command_buffer
+                           withRange:NSMakeRange(body->location, body->length)];
+      break;
+    }
     }
 
     next = next->next.ptr;
@@ -841,6 +865,12 @@ _MTLComputeCommandEncoder_encodeCommands(void *obj) {
       struct wmtcmd_compute_memory_barrier *body = (struct wmtcmd_compute_memory_barrier *)next;
       [encoder memoryBarrierWithScope:(MTLBarrierScope)body->scope];
       break;
+    }    
+    case WMTComputeCommandExecuteCommandsInBuffer: {
+      struct wmtcmd_compute_executecommands *body = (struct wmtcmd_compute_executecommands *)next;
+      [encoder executeCommandsInBuffer:(id<MTLIndirectCommandBuffer>)body->indirect_command_buffer
+                             withRange:NSMakeRange(body->location, body->length)];
+      break;
     }
     }
 
@@ -954,6 +984,16 @@ _MTLRenderCommandEncoder_encodeCommands(void *obj) {
       [encoder setStencilReferenceValue:body->stencil_ref];
       break;
     }
+    case WMTRenderCommandSetBlendFactor: {
+      struct wmtcmd_render_setblendcolor *body = (struct wmtcmd_render_setblendcolor *)next;
+      [encoder setBlendColorRed:body->red green:body->green blue:body->blue alpha:body->alpha];
+      break;
+    }
+    case WMTRenderCommandSetStencilRef: {
+      struct wmtcmd_render_setstencilref *body = (struct wmtcmd_render_setstencilref *)next;
+      [encoder setStencilReferenceValue:body->stencil_ref];
+      break;
+    }
     case WMTRenderCommandSetVisibilityMode: {
       struct wmtcmd_render_setvisibilitymode *body = (struct wmtcmd_render_setvisibilitymode *)next;
       [encoder setVisibilityResultMode:(MTLVisibilityResultMode)body->mode offset:body->offset];
@@ -1174,6 +1214,12 @@ _MTLRenderCommandEncoder_encodeCommands(void *obj) {
       [encoder dispatchThreadsPerTile:MTLSizeMake(body->width, body->height, 1)];
       break;
     }
+    case WMTRenderCommandExecuteCommandsInBuffer: {
+      struct wmtcmd_render_executecommands *body = (struct wmtcmd_render_executecommands *)next;
+      [encoder executeCommandsInBuffer:(id<MTLIndirectCommandBuffer>)body->indirect_command_buffer
+                             withRange:NSMakeRange(body->location, body->length)];
+      break;
+    }
     }
     next = next->next.ptr;
   }
@@ -2908,6 +2954,144 @@ _MTLCommandQueue_addResidencySet(void *obj) {
   return STATUS_SUCCESS;
 }
 
+static NTSTATUS
+_MTLDevice_newHeap(void *obj) {
+  struct unixcall_mtldevice_newheap *params = obj;
+  id<MTLDevice> device = (id<MTLDevice>)params->device;
+  struct WMTHeapInfo const *info = params->info.ptr;
+  MTLHeapDescriptor *desc = [[MTLHeapDescriptor alloc] init];
+  desc.resourceOptions = (MTLResourceOptions)info->options;
+  desc.sparsePageSize = (MTLSparsePageSize)info->sparse_page_size;
+  desc.type = (MTLHeapType)info->type;
+  desc.size = info->size;
+
+  params->ret = (obj_handle_t)[device newHeapWithDescriptor:desc];
+
+  [desc release];
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLDevice_heapBufferSizeAndAlign(void *obj) {
+  struct unixcall_mtldevice_heapbuffersizeandalign *params = obj;
+  id<MTLDevice> device = (id<MTLDevice>)params->device;
+  MTLSizeAndAlign ret =
+      [device heapBufferSizeAndAlignWithLength:params->length options:(MTLResourceOptions)params->options];
+  params->ret_size = ret.size;
+  params->ret_align = ret.align;
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLDevice_heapTextureSizeAndAlign(void *obj) {
+  struct unixcall_mtldevice_heaptexturesizeandalign *params = obj;
+  id<MTLDevice> device = (id<MTLDevice>)params->device;
+  struct WMTTextureInfo *info = params->info.ptr;
+  MTLTextureDescriptor *desc = [[MTLTextureDescriptor alloc] init];
+  fill_texture_descriptor(desc, info);
+  MTLSizeAndAlign ret = [device heapTextureSizeAndAlignWithDescriptor:desc];
+  params->ret_size = ret.size;
+  params->ret_align = ret.align;
+  [desc release];
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLHeap_newBuffer(void *obj) {
+  struct unixcall_mtlheap_newbuffer *params = obj;
+  id<MTLHeap> heap = (id<MTLHeap>)params->heap;
+  struct WMTBufferInfo *info = params->info.ptr;
+  id<MTLBuffer> buffer;
+  if (~params->offset == 0)
+    buffer = [heap newBufferWithLength:info->length options:(enum MTLResourceOptions)info->options];
+  else
+    buffer =
+        [heap newBufferWithLength:info->length options:(enum MTLResourceOptions)info->options offset:params->offset];
+  info->memory.ptr = [heap storageMode] == MTLStorageModePrivate ? NULL : [buffer contents];
+
+  params->ret = (obj_handle_t)buffer;
+  info->gpu_address = [buffer gpuAddress];
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLHeap_newTexture(void *obj) {
+  struct unixcall_mtlheap_newtexture *params = obj;
+  id<MTLHeap> heap = (id<MTLHeap>)params->heap;
+  struct WMTTextureInfo *info = params->info.ptr;
+  MTLTextureDescriptor *desc = [[MTLTextureDescriptor alloc] init];
+  fill_texture_descriptor(desc, info);
+
+  id<MTLTexture> ret;
+  if (~params->offset == 0)
+    ret = [heap newTextureWithDescriptor:desc];
+  else
+    ret = [heap newTextureWithDescriptor:desc offset:params->offset];
+  params->ret = (obj_handle_t)ret;
+  info->gpu_resource_id = [ret gpuResourceID]._impl;
+  info->mach_port = 0;
+
+  [desc release];
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLDevice_newIndirectCommandBuffer(void *obj) {
+  struct unixcall_mtldevice_newicb *params = obj;
+  id<MTLDevice> device = (id<MTLDevice>)params->device;
+
+  struct WMTIndirectCommandBufferInfo *info = params->info.ptr;
+
+  MTLIndirectCommandBufferDescriptor *desc = [[MTLIndirectCommandBufferDescriptor alloc] init];
+
+  desc.commandTypes = (MTLIndirectCommandType)info->type;
+
+  desc.maxFragmentBufferBindCount = info->max_fragment_buffer_binding;
+  desc.maxKernelBufferBindCount = info->max_kernel_buffer_binding;
+  desc.maxKernelThreadgroupMemoryBindCount = info->max_kernel_threadgroup_memory_binding;
+  desc.maxMeshBufferBindCount = info->max_mesh_buffer_binding;
+  desc.maxObjectBufferBindCount = info->max_object_buffer_binding;
+  desc.maxObjectThreadgroupMemoryBindCount = info->max_object_threadgroup_memory_binding;
+  desc.maxVertexBufferBindCount = info->max_vertex_buffer_binding;
+
+  desc.inheritBuffers = info->inherit_buffers;
+  desc.inheritPipelineState = info->inherit_pso;
+  desc.supportDynamicAttributeStride = info->support_dynamic_attribute_stride;
+  desc.supportRayTracing = info->support_ray_tracing;
+
+#if 0 // needs macOS 26 SDK, and we don't really use it at the moment
+  desc.inheritCullMode = info->inherit_cull_mode;
+  desc.inheritDepthBias = info->inherit_depth_bias;
+  desc.inheritDepthClipMode = info->inherit_depth_clip_mode;
+  desc.inheritDepthStencilState = info->inherit_depth_stencil_state;
+  desc.inheritFrontFacingWinding = info->inherit_front_facing;
+  desc.inheritTriangleFillMode = info->inherit_fill_mode;
+  desc.supportColorAttachmentMapping = info->support_color_attachment_mapping;
+#endif
+
+  id<MTLIndirectCommandBuffer> icb =
+      [device newIndirectCommandBufferWithDescriptor:desc
+                                     maxCommandCount:params->max_count
+                                             options:(MTLResourceOptions)params->options];
+  params->ret = (obj_handle_t)icb;
+  info->gpu_resource_id = [icb gpuResourceID]._impl;
+  [desc release];
+  return STATUS_SUCCESS;
+}
+
+static NTSTATUS
+_MTLDevice_newLibraryWithSource(void *obj) {
+  struct unixcall_mtldevice_newlibrarywithsource *params = obj;
+  id<MTLDevice> device = (id<MTLDevice>)params->device;
+  NSString *source =
+      [[NSString alloc] initWithBytes:params->source.ptr length:params->length encoding:NSASCIIStringEncoding];
+  NSError *err = NULL;
+  params->ret_library = (obj_handle_t)[device newLibraryWithSource:source options:nil error:&err];
+  params->ret_error = (obj_handle_t)err;
+  [source release];
+  return STATUS_SUCCESS;
+}
+
 /*
  * Definition from cache.c
  */
@@ -3057,6 +3241,13 @@ const void *__wine_unix_call_funcs[] = {
     &_MTLResidencySet_removeAllAllocations,
     &_MTLResidencySet_commit,
     &_MTLCommandQueue_addResidencySet,
+    &_MTLDevice_newHeap,
+    &_MTLDevice_heapBufferSizeAndAlign,
+    &_MTLDevice_heapTextureSizeAndAlign,
+    &_MTLHeap_newBuffer,
+    &_MTLHeap_newTexture,
+    &_MTLDevice_newIndirectCommandBuffer,
+    &_MTLDevice_newLibraryWithSource,
 };
 
 #ifndef DXMT_NATIVE
@@ -3199,5 +3390,12 @@ const void *__wine_unix_call_wow64_funcs[] = {
     &_MTLResidencySet_removeAllAllocations,
     &_MTLResidencySet_commit,
     &_MTLCommandQueue_addResidencySet,
+    &_MTLDevice_newHeap,
+    &_MTLDevice_heapBufferSizeAndAlign,
+    &_MTLDevice_heapTextureSizeAndAlign,
+    &_MTLHeap_newBuffer,
+    &_MTLHeap_newTexture,
+    &_MTLDevice_newIndirectCommandBuffer,
+    &_MTLDevice_newLibraryWithSource,
 };
 #endif
diff --git a/src/winemetal/winemetal.h b/src/winemetal/winemetal.h
index 44296d041..a72cc8eb9 100644
--- a/src/winemetal/winemetal.h
+++ b/src/winemetal/winemetal.h
@@ -616,7 +616,7 @@ struct WMTComputePipelineInfo {
   obj_handle_t binary_archive_for_serialization;
   uint8_t num_binary_archives_for_lookup;
   bool fail_on_binary_archive_miss;
-  uint8_t padding;
+  bool support_indirect_command_buffers;
   bool tgsize_is_multiple_of_sgwidth;
   uint32_t immutable_buffers;
 };
@@ -827,7 +827,8 @@ struct WMTRenderPipelineInfo {
   struct WMTConstMemoryPointer binary_archives_for_lookup;
   uint8_t num_binary_archives_for_lookup;
   bool fail_on_binary_archive_miss;
-  uint8_t padding[6];
+  bool support_indirect_command_buffers;
+  uint8_t padding[5];
 };
 
 struct WMTMeshRenderPipelineInfo {
@@ -852,7 +853,8 @@ struct WMTMeshRenderPipelineInfo {
   struct WMTConstMemoryPointer binary_archives_for_lookup;
   uint8_t num_binary_archives_for_lookup;
   bool fail_on_binary_archive_miss;
-  uint8_t padding[6];
+  bool support_indirect_command_buffers;
+  uint8_t padding[5];
 };
 
 WINEMETAL_API obj_handle_t
@@ -886,6 +888,8 @@ enum WMTBlitCommandType : uint16_t {
   WMTBlitCommandFillBuffer,
   WMTBlitCommandResolveCounters,
   WMTBlitCommandCopyFromBufferToTextureWithBlitOption,
+  WMTBlitCommandCopyFromTextureToBufferWithBlitOption,
+  WMTBlitCommandResetCommandsInBuffer,
 };
 
 enum WMTBlitOption : uint16_t {
@@ -979,6 +983,22 @@ struct wmtcmd_blit_copy_from_texture_to_buffer {
   uint32_t bytes_per_image;
 };
 
+struct wmtcmd_blit_copy_from_texture_to_buffer_withblitoption {
+  enum WMTBlitCommandType type;
+  uint16_t reserved[3];
+  struct WMTMemoryPointer next;
+  obj_handle_t src;
+  uint32_t slice;
+  uint16_t level;
+  uint16_t options;
+  struct WMTOrigin origin;
+  struct WMTSize size;
+  obj_handle_t dst;
+  uint64_t offset;
+  uint32_t bytes_per_row;
+  uint32_t bytes_per_image;
+};
+
 struct wmtcmd_blit_generate_mipmaps {
   enum WMTBlitCommandType type;
   uint16_t reserved[3];
@@ -1014,6 +1034,15 @@ struct wmtcmd_blit_resolvecounters {
   uint64_t dst_offset;
 };
 
+struct wmtcmd_blit_resetcommands {
+  enum WMTBlitCommandType type;
+  uint16_t reserved[3];
+  struct WMTMemoryPointer next;
+  obj_handle_t indirect_command_buffer;
+  uint32_t location;
+  uint32_t length;
+};
+
 WINEMETAL_API void MTLBlitCommandEncoder_encodeCommands(obj_handle_t encoder, const struct wmtcmd_base *cmd_head);
 
 enum WMTComputeCommandType : uint16_t {
@@ -1030,6 +1059,7 @@ enum WMTComputeCommandType : uint16_t {
   WMTComputeCommandWaitForFence,
   WMTComputeCommandUpdateFence,
   WMTComputeCommandMemoryBarrier,
+  WMTComputeCommandExecuteCommandsInBuffer,
 };
 
 struct wmtcmd_compute_nop {
@@ -1129,6 +1159,15 @@ enum WMTComputeCommandType type;
   enum WMTBarrierScope scope;
 };
 
+struct wmtcmd_compute_executecommands {
+  enum WMTComputeCommandType type;
+  uint16_t reserved[3];
+  struct WMTMemoryPointer next;
+  obj_handle_t indirect_command_buffer;
+  uint32_t location;
+  uint32_t length;
+};
+
 WINEMETAL_API void MTLComputeCommandEncoder_encodeCommands(obj_handle_t encoder, const struct wmtcmd_base *cmd_head);
 
 enum WMTRenderCommandType : uint16_t {
@@ -1174,6 +1213,9 @@ enum WMTRenderCommandType : uint16_t {
   WMTRenderCommandDXMTTessellationMeshDrawIndirect,
   WMTRenderCommandDXMTTessellationMeshDrawIndexedIndirect,
   WMTRenderCommandDispatchThreadsPerTile,
+  WMTRenderCommandExecuteCommandsInBuffer,
+  WMTRenderCommandSetBlendFactor,
+  WMTRenderCommandSetStencilRef,
 };
 
 struct wmtcmd_render_nop {
@@ -1435,6 +1477,13 @@ struct wmtcmd_render_setblendcolor {
   uint8_t stencil_ref;
 };
 
+struct wmtcmd_render_setstencilref {
+  enum WMTRenderCommandType type;
+  uint16_t reserved[3];
+  struct WMTMemoryPointer next;
+  uint8_t stencil_ref;
+};
+
 struct wmtcmd_render_dxmt_geometry_draw {
   enum WMTRenderCommandType type;
   uint16_t reserved[3];
@@ -1551,6 +1600,15 @@ struct wmtcmd_render_dispatch_threads_per_tile {
   uint32_t height;
 };
 
+struct wmtcmd_render_executecommands {
+  enum WMTRenderCommandType type;
+  uint16_t reserved[3];
+  struct WMTMemoryPointer next;
+  obj_handle_t indirect_command_buffer;
+  uint32_t location;
+  uint32_t length;
+};
+
 WINEMETAL_API void MTLRenderCommandEncoder_encodeCommands(obj_handle_t encoder, const struct wmtcmd_base *cmd_head);
 
 WINEMETAL_API enum WMTPixelFormat MTLTexture_pixelFormat(obj_handle_t texture);
@@ -1985,4 +2043,84 @@ WINEMETAL_API void MTLResidencySet_commit(obj_handle_t residency_set);
 
 WINEMETAL_API void MTLCommandQueue_addResidencySet(obj_handle_t queue, obj_handle_t residency_set);
 
+enum WMTHeapType : uint32_t {
+  WMTHeapTypeAutomatic = 0,
+  WMTHeapTypePlacement = 1,
+  WMTHeapTypeSparse = 2,
+};
+enum WMTSparsePageSize : uint32_t {
+  WMTSparsePageSize16 = 101,
+  WMTSparsePageSize64 = 102,
+  WMTSparsePageSize256 = 103,
+};
+struct WMTHeapInfo {
+  uint64_t size;
+  enum WMTResourceOptions options;
+  enum WMTHeapType type;
+  enum WMTSparsePageSize sparse_page_size;
+};
+
+STATIC_ASSERT(sizeof(WMTHeapInfo) == 24);
+
+WINEMETAL_API obj_handle_t MTLDevice_newHeap(obj_handle_t device, const struct WMTHeapInfo *info);
+
+struct WMTSizeAndAlign {
+  // 32-bit is sufficient
+  uint32_t size;
+  uint32_t align;
+};
+
+WINEMETAL_API struct WMTSizeAndAlign
+MTLDevice_heapBufferSizeAndAlign(obj_handle_t device, uint64_t length, enum WMTResourceOptions options);
+
+WINEMETAL_API struct WMTSizeAndAlign
+MTLDevice_heapTextureSizeAndAlign(obj_handle_t device, struct WMTTextureInfo *info);
+
+WINEMETAL_API obj_handle_t MTLHeap_newBuffer(obj_handle_t heap, struct WMTBufferInfo *info, uint64_t offset);
+
+WINEMETAL_API obj_handle_t MTLHeap_newTexture(obj_handle_t heap, struct WMTTextureInfo *info, uint64_t offset);
+
+enum WMTIndirectCommandType {
+  WMTIndirectCommandTypeDraw = (1 << 0),
+  WMTIndirectCommandTypeDrawIndexed = (1 << 1),
+  WMTIndirectCommandTypeDrawPatches = (1 << 2),
+  WMTIndirectCommandTypeDrawIndexedPatches = (1 << 3),
+  WMTIndirectCommandTypeConcurrentDispatch = (1 << 5),
+  WMTIndirectCommandTypeConcurrentDispatchThreads = (1 << 6),
+  WMTIndirectCommandTypeDrawMeshThreadgroups = (1 << 7),
+  WMTIndirectCommandTypeDrawMeshThreads = (1 << 8),
+};
+
+struct WMTIndirectCommandBufferInfo {
+  enum WMTIndirectCommandType type;
+  uint32_t inherit_buffers                  : 1;
+  uint32_t inherit_pso                      : 1;
+  uint32_t inherit_cull_mode                : 1;
+  uint32_t inherit_depth_bias               : 1;
+  uint32_t inherit_depth_clip_mode          : 1;
+  uint32_t inherit_depth_stencil_state      : 1;
+  uint32_t inherit_front_facing             : 1;
+  uint32_t inherit_fill_mode                : 1;
+  uint32_t support_ray_tracing              : 1;
+  uint32_t support_dynamic_attribute_stride : 1;
+  uint32_t support_color_attachment_mapping : 1;
+  uint32_t padding_bits                     : 21;
+  uint8_t max_vertex_buffer_binding;
+  uint8_t max_fragment_buffer_binding;
+  uint8_t max_kernel_buffer_binding;
+  uint8_t max_mesh_buffer_binding;
+  uint8_t max_object_buffer_binding;
+  uint8_t max_object_threadgroup_memory_binding;
+  uint8_t max_kernel_threadgroup_memory_binding;
+  uint8_t padding;
+  uint64_t gpu_resource_id; // out
+};
+
+WINEMETAL_API obj_handle_t MTLDevice_newIndirectCommandBuffer(
+    obj_handle_t device, struct WMTIndirectCommandBufferInfo *info, uint64_t max_count, enum WMTResourceOptions options
+);
+
+WINEMETAL_API obj_handle_t
+MTLDevice_newLibraryWithSource(obj_handle_t device, const char *source, uint64_t length, obj_handle_t *err_out);
+
 #endif
\ No newline at end of file
diff --git a/src/winemetal/winemetal_thunks.c b/src/winemetal/winemetal_thunks.c
index a4da50a9d..f20c774af 100644
--- a/src/winemetal/winemetal_thunks.c
+++ b/src/winemetal/winemetal_thunks.c
@@ -1232,3 +1232,84 @@ MTLCommandQueue_addResidencySet(obj_handle_t queue, obj_handle_t residency_set)
   params.arg = residency_set;
   UNIX_CALL(137, &params);
 }
+
+WINEMETAL_API obj_handle_t
+MTLDevice_newHeap(obj_handle_t device, const struct WMTHeapInfo *info) {
+  struct unixcall_mtldevice_newheap params;
+  params.device = device;
+  WMT_MEMPTR_SET(params.info, info);
+  UNIX_CALL(138, &params);
+  return params.ret;
+}
+
+WINEMETAL_API struct WMTSizeAndAlign
+MTLDevice_heapBufferSizeAndAlign(obj_handle_t device, uint64_t length, enum WMTResourceOptions options) {
+  struct unixcall_mtldevice_heapbuffersizeandalign params;
+  params.device = device;
+  params.length = length;
+  params.options = options;
+  UNIX_CALL(139, &params);
+  struct WMTSizeAndAlign ret;
+  ret.size = params.ret_size;
+  ret.align = params.ret_align;
+  return ret;
+}
+
+WINEMETAL_API struct WMTSizeAndAlign
+MTLDevice_heapTextureSizeAndAlign(obj_handle_t device, struct WMTTextureInfo *info) {
+  struct unixcall_mtldevice_heaptexturesizeandalign params;
+  params.device = device;
+  WMT_MEMPTR_SET(params.info, info);
+  UNIX_CALL(140, &params);
+  struct WMTSizeAndAlign ret;
+  ret.size = params.ret_size;
+  ret.align = params.ret_align;
+  return ret;
+}
+
+WINEMETAL_API obj_handle_t
+MTLHeap_newBuffer(obj_handle_t heap, struct WMTBufferInfo *info, uint64_t offset) {
+  struct unixcall_mtlheap_newbuffer params;
+  params.heap = heap;
+  WMT_MEMPTR_SET(params.info, info);
+  params.offset = offset;
+  UNIX_CALL(141, &params);
+  return params.ret;
+}
+
+WINEMETAL_API obj_handle_t
+MTLHeap_newTexture(obj_handle_t heap, struct WMTTextureInfo *info, uint64_t offset) {
+  struct unixcall_mtlheap_newtexture params;
+  params.heap = heap;
+  WMT_MEMPTR_SET(params.info, info);
+  params.offset = offset;
+  UNIX_CALL(142, &params);
+  return params.ret;
+}
+
+WINEMETAL_API obj_handle_t
+MTLDevice_newIndirectCommandBuffer(
+    obj_handle_t device, struct WMTIndirectCommandBufferInfo *info, uint64_t max_count, enum WMTResourceOptions options
+) {
+  struct unixcall_mtldevice_newicb params;
+  params.device = device;
+  WMT_MEMPTR_SET(params.info, info);
+  params.max_count = max_count;
+  params.options = options;
+  UNIX_CALL(143, &params);
+  return params.ret;
+}
+
+WINEMETAL_API obj_handle_t
+MTLDevice_newLibraryWithSource(obj_handle_t device, const char *source, uint64_t length, obj_handle_t *err_out) {
+  struct unixcall_mtldevice_newlibrarywithsource params;
+  params.device = device;
+  WMT_MEMPTR_SET(params.source, source);
+  params.length = length;
+  params.ret_error = 0;
+  params.ret_library = 0;
+  UNIX_CALL(144, &params);
+  if (err_out)
+    *err_out = params.ret_error;
+  return params.ret_library;
+}
diff --git a/src/winemetal/winemetal_thunks.h b/src/winemetal/winemetal_thunks.h
index 093a49d58..37645c959 100644
--- a/src/winemetal/winemetal_thunks.h
+++ b/src/winemetal/winemetal_thunks.h
@@ -117,6 +117,14 @@ struct unixcall_mtldevice_newlibrary {
   obj_handle_t ret_library;
 };
 
+struct unixcall_mtldevice_newlibrarywithsource {
+  obj_handle_t device;
+  struct WMTConstMemoryPointer source;
+  uint64_t length;
+  obj_handle_t ret_error;
+  obj_handle_t ret_library;
+};
+
 struct unixcall_mtldevice_newcomputepso {
   obj_handle_t device;
   struct WMTConstMemoryPointer info;
@@ -384,6 +392,49 @@ struct unixcall_mtlresidencyset_addallocations {
   uint64_t count;
 };
 
+struct unixcall_mtldevice_newheap {
+  obj_handle_t device;
+  struct WMTConstMemoryPointer info;
+  obj_handle_t ret;
+};
+
+struct unixcall_mtldevice_heapbuffersizeandalign {
+  obj_handle_t device;
+  uint64_t length;
+  enum WMTResourceOptions options;
+  uint32_t ret_size;
+  uint32_t ret_align;
+};
+
+struct unixcall_mtldevice_heaptexturesizeandalign {
+  obj_handle_t device;
+  struct WMTMemoryPointer info;
+  uint32_t ret_size;
+  uint32_t ret_align;
+};
+
+struct unixcall_mtlheap_newbuffer {
+  obj_handle_t heap;
+  struct WMTMemoryPointer info;
+  uint64_t offset;
+  obj_handle_t ret;
+};
+
+struct unixcall_mtlheap_newtexture {
+  obj_handle_t heap;
+  struct WMTMemoryPointer info;
+  uint64_t offset;
+  obj_handle_t ret;
+};
+
+struct unixcall_mtldevice_newicb {
+  obj_handle_t device;
+  struct WMTMemoryPointer info;
+  uint64_t max_count;
+  uint64_t options;
+  obj_handle_t ret;
+};
+
 #pragma pack(pop)
 
 #endif