diff --git a/src/kosmickrisp/bridge/mtl_encoder.h b/src/kosmickrisp/bridge/mtl_encoder.h
index bfd582d571e..3daac982196 100644
--- a/src/kosmickrisp/bridge/mtl_encoder.h
+++ b/src/kosmickrisp/bridge/mtl_encoder.h
@@ -50,13 +50,6 @@ void mtl_compute_set_pipeline_state(mtl_compute_encoder *encoder,
 void mtl_compute_set_buffer(mtl_compute_encoder *encoder, mtl_buffer *buffer,
                             size_t offset, size_t index);
 
-void mtl_compute_use_resource(mtl_compute_encoder *encoder,
-                              mtl_resource *res_handle, uint32_t usage);
-
-void mtl_compute_use_resources(mtl_compute_encoder *encoder,
-                               mtl_resource **resource_handles, uint32_t count,
-                               enum mtl_resource_usage usage);
-
 void mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
                            uint32_t count);
 
diff --git a/src/kosmickrisp/bridge/mtl_encoder.m b/src/kosmickrisp/bridge/mtl_encoder.m
index 8385ea8ac0b..3f88c547e6d 100644
--- a/src/kosmickrisp/bridge/mtl_encoder.m
+++ b/src/kosmickrisp/bridge/mtl_encoder.m
@@ -191,29 +191,6 @@ mtl_compute_set_buffer(mtl_compute_encoder *encoder,
    }
 }
 
-void
-mtl_compute_use_resource(mtl_compute_encoder *encoder,
-                         mtl_resource *res_handle, uint32_t usage)
-{
-   @autoreleasepool {
-      id<MTLComputeCommandEncoder> enc = (id<MTLComputeCommandEncoder>)encoder;
-      id<MTLResource> res = (id<MTLResource>)res_handle;
-      [enc useResource:res usage:(MTLResourceUsage)usage];
-   }
-}
-
-void
-mtl_compute_use_resources(mtl_compute_encoder *encoder,
-                          mtl_resource **resource_handles, uint32_t count,
-                          enum mtl_resource_usage usage)
-{
-   @autoreleasepool {
-      id<MTLComputeCommandEncoder> enc = (id<MTLComputeCommandEncoder>)encoder;
-      id<MTLResource> *handles = (id<MTLResource>*)resource_handles;
-      [enc useResources:handles count:count usage:(MTLResourceUsage)usage];
-   }
-}
-
 void
 mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
                       uint32_t count)
diff --git a/src/kosmickrisp/bridge/stubs/mtl_encoder.c b/src/kosmickrisp/bridge/stubs/mtl_encoder.c
index 672e628c1dc..d2886d5059e 100644
--- a/src/kosmickrisp/bridge/stubs/mtl_encoder.c
+++ b/src/kosmickrisp/bridge/stubs/mtl_encoder.c
@@ -88,19 +88,6 @@ mtl_compute_set_buffer(mtl_compute_encoder *encoder, mtl_buffer *buffer,
 {
 }
 
-void
-mtl_compute_use_resource(mtl_compute_encoder *encoder, mtl_resource *res_handle,
-                         uint32_t usage)
-{
-}
-
-void
-mtl_compute_use_resources(mtl_compute_encoder *encoder,
-                          mtl_resource **resource_handles, uint32_t count,
-                          enum mtl_resource_usage usage)
-{
-}
-
 void
 mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
                       uint32_t count)
diff --git a/src/kosmickrisp/vulkan/kk_bo.c b/src/kosmickrisp/vulkan/kk_bo.c
index 708b93fb2ed..eb9b507b86f 100644
--- a/src/kosmickrisp/vulkan/kk_bo.c
+++ b/src/kosmickrisp/vulkan/kk_bo.c
@@ -50,6 +50,8 @@ kk_alloc_bo(struct kk_device *dev, struct vk_object_base *log_obj,
    bo->gpu = mtl_buffer_get_gpu_address(map);
    bo->cpu = mtl_get_contents(map);
 
+   kk_device_add_heap_to_residency_set(dev, handle);
+
    *bo_out = bo;
    return result;
 
@@ -64,6 +66,7 @@ fail_heap:
 void
 kk_destroy_bo(struct kk_device *dev, struct kk_bo *bo)
 {
+   kk_device_remove_heap_from_residency_set(dev, bo->mtl_handle);
    mtl_release(bo->map);
    mtl_release(bo->mtl_handle);
    FREE(bo);
diff --git a/src/kosmickrisp/vulkan/kk_cmd_buffer.c b/src/kosmickrisp/vulkan/kk_cmd_buffer.c
index b38b2f2e844..d438f7f0fe8 100644
--- a/src/kosmickrisp/vulkan/kk_cmd_buffer.c
+++ b/src/kosmickrisp/vulkan/kk_cmd_buffer.c
@@ -29,7 +29,6 @@ kk_descriptor_state_fini(struct kk_cmd_buffer *cmd,
       desc->push[i] = NULL;
       desc->sets[i] = NULL; /* We also need to set sets to NULL so state doesn't
                                propagate if we reset it */
-      desc->sets_not_resident = 0u;
    }
 }
 
@@ -204,8 +203,6 @@ kk_bind_descriptor_sets(struct kk_descriptor_state *desc,
          }
          desc->sets[s] = set;
 
-         desc->sets_not_resident |= BITFIELD_BIT(s);
-
          /* Binding descriptors invalidates push descriptors */
          desc->push_dirty &= ~BITFIELD_BIT(s);
       }
@@ -262,11 +259,9 @@ kk_cmd_push_descriptors(struct kk_cmd_buffer *cmd,
                         struct kk_descriptor_set_layout *set_layout,
                         uint32_t set)
 {
-   struct kk_device *dev = kk_cmd_buffer_device(cmd);
    assert(set < KK_MAX_SETS);
    if (unlikely(desc->push[set] == NULL)) {
-      size_t size = sizeof(*desc->push[set]) +
-                    (sizeof(mtl_resource *) * set_layout->descriptor_count);
+      size_t size = sizeof(*desc->push[set]);
       desc->push[set] = vk_zalloc(&cmd->vk.pool->alloc, size, 8,
                                   VK_SYSTEM_ALLOCATION_SCOPE_OBJECT);
       if (unlikely(desc->push[set] == NULL)) {
@@ -274,14 +269,11 @@ kk_cmd_push_descriptors(struct kk_cmd_buffer *cmd,
          return NULL;
       }
       desc->push[set]->layout = set_layout;
-      for (uint32_t i = 0u; i < set_layout->descriptor_count; ++i)
-         desc->push[set]->mtl_resources[i] = dev->null_descriptor->map;
    }
 
    /* Pushing descriptors replaces whatever sets are bound */
    desc->sets[set] = NULL;
    desc->push_dirty |= BITFIELD_BIT(set);
-   desc->sets_not_resident |= BITFIELD_BIT(set);
 
    return desc->push[set];
 }
@@ -425,7 +417,6 @@ kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
          return;
 
       memcpy(bo->cpu, push_set->data, sizeof(push_set->data));
-      push_set->mtl_descriptor_buffer = bo->map;
       desc->root.sets[set_idx] = bo->gpu;
       desc->set_sizes[set_idx] = sizeof(push_set->data);
    }
@@ -434,77 +425,12 @@ kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
    desc->push_dirty = 0;
 }
 
-static void
-kk_make_graphics_descriptor_resources_resident(struct kk_cmd_buffer *cmd)
-{
-   struct kk_descriptor_state *desc = &cmd->state.gfx.descriptors;
-   mtl_render_encoder *encoder = kk_render_encoder(cmd);
-   /* Make resources resident as required by Metal */
-   u_foreach_bit(set_index, desc->sets_not_resident) {
-      mtl_resource *descriptor_buffer = NULL;
-
-      /* If we have no set, it means it was a push set */
-      if (desc->sets[set_index]) {
-         struct kk_descriptor_set *set = desc->sets[set_index];
-         descriptor_buffer = set->mtl_descriptor_buffer;
-      } else {
-         struct kk_push_descriptor_set *push_set = desc->push[set_index];
-         descriptor_buffer = push_set->mtl_descriptor_buffer;
-      }
-
-      /* We could have empty descriptor sets for some reason... */
-      if (descriptor_buffer) {
-         mtl_render_use_resource(encoder, descriptor_buffer,
-                                 MTL_RESOURCE_USAGE_READ);
-      }
-   }
-
-   desc->sets_not_resident = 0u;
-}
-
-static void
-kk_make_compute_descriptor_resources_resident(struct kk_cmd_buffer *cmd)
-{
-   struct kk_descriptor_state *desc = &cmd->state.cs.descriptors;
-   mtl_compute_encoder *encoder = kk_compute_encoder(cmd);
-   u_foreach_bit(set_index, desc->sets_not_resident) {
-      /* Make resources resident as required by Metal */
-      mtl_resource *descriptor_buffer = NULL;
-      if (desc->sets[set_index]) {
-         struct kk_descriptor_set *set = desc->sets[set_index];
-         descriptor_buffer = set->mtl_descriptor_buffer;
-      } else {
-         struct kk_push_descriptor_set *push_set = desc->push[set_index];
-         descriptor_buffer = push_set->mtl_descriptor_buffer;
-      }
-
-      /* We could have empty descriptor sets for some reason... */
-      if (descriptor_buffer) {
-         mtl_compute_use_resource(encoder, descriptor_buffer,
-                                  MTL_RESOURCE_USAGE_READ);
-      }
-   }
-
-   desc->sets_not_resident = 0u;
-}
-
-void
-kk_make_descriptor_resources_resident(struct kk_cmd_buffer *cmd,
-                                      VkPipelineBindPoint bind_point)
-{
-   if (bind_point == VK_PIPELINE_BIND_POINT_GRAPHICS)
-      kk_make_graphics_descriptor_resources_resident(cmd);
-   else if (bind_point == VK_PIPELINE_BIND_POINT_COMPUTE)
-      kk_make_compute_descriptor_resources_resident(cmd);
-}
-
 void
 kk_cmd_write(struct kk_cmd_buffer *cmd, mtl_buffer *buffer, uint64_t addr,
              uint64_t value)
 {
    util_dynarray_append(&cmd->encoder->imm_writes, addr);
    util_dynarray_append(&cmd->encoder->imm_writes, value);
-   util_dynarray_append(&cmd->encoder->resident_buffers, buffer);
 }
 
 VKAPI_ATTR void VKAPI_CALL
diff --git a/src/kosmickrisp/vulkan/kk_cmd_buffer.h b/src/kosmickrisp/vulkan/kk_cmd_buffer.h
index d5722a53cbb..43f2a8e2e50 100644
--- a/src/kosmickrisp/vulkan/kk_cmd_buffer.h
+++ b/src/kosmickrisp/vulkan/kk_cmd_buffer.h
@@ -61,10 +61,6 @@ struct kk_descriptor_state {
 
    uint32_t set_sizes[KK_MAX_SETS];
    struct kk_descriptor_set *sets[KK_MAX_SETS];
-   mtl_resource **resources[KK_MAX_SETS];
-   /* Non resident sets can either be sets or push. If sets[index] == NULL, then
-    * push[index] != NULL */
-   uint32_t sets_not_resident;
 
    uint32_t push_dirty;
    struct kk_push_descriptor_set *push[KK_MAX_SETS];
@@ -255,9 +251,6 @@ uint64_t kk_upload_descriptor_root(struct kk_cmd_buffer *cmd,
 void kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
                                           struct kk_descriptor_state *desc);
 
-void kk_make_descriptor_resources_resident(struct kk_cmd_buffer *cmd,
-                                           VkPipelineBindPoint bind_point);
-
 void kk_cmd_write(struct kk_cmd_buffer *cmd, mtl_buffer *buffer, uint64_t addr,
                   uint64_t value);
 
diff --git a/src/kosmickrisp/vulkan/kk_cmd_dispatch.c b/src/kosmickrisp/vulkan/kk_cmd_dispatch.c
index bfca701c52f..2a66ad47794 100644
--- a/src/kosmickrisp/vulkan/kk_cmd_dispatch.c
+++ b/src/kosmickrisp/vulkan/kk_cmd_dispatch.c
@@ -72,18 +72,12 @@ static void
 kk_flush_compute_state(struct kk_cmd_buffer *cmd)
 {
    mtl_compute_encoder *enc = kk_compute_encoder(cmd);
-   struct kk_device *dev = kk_cmd_buffer_device(cmd);
 
    // Fill Metal argument buffer with descriptor set addresses
    struct kk_descriptor_state *desc = &cmd->state.cs.descriptors;
 
    if (desc->push_dirty)
       kk_cmd_buffer_flush_push_descriptors(cmd, desc);
-   /* After push descriptors' buffers are created. Otherwise, the buffer where
-    * they live will not be created and cannot make it resident */
-   if (desc->sets_not_resident)
-      kk_make_descriptor_resources_resident(cmd,
-                                            VK_PIPELINE_BIND_POINT_COMPUTE);
    if (desc->root_dirty)
       kk_upload_descriptor_root(cmd, VK_PIPELINE_BIND_POINT_COMPUTE);
 
diff --git a/src/kosmickrisp/vulkan/kk_cmd_draw.c b/src/kosmickrisp/vulkan/kk_cmd_draw.c
index 3a9b69a1caa..7bf2ee4c1f9 100644
--- a/src/kosmickrisp/vulkan/kk_cmd_draw.c
+++ b/src/kosmickrisp/vulkan/kk_cmd_draw.c
@@ -816,11 +816,6 @@ kk_flush_draw_state(struct kk_cmd_buffer *cmd)
 
    if (desc->push_dirty)
       kk_cmd_buffer_flush_push_descriptors(cmd, desc);
-   /* After push descriptors' buffers are created. Otherwise, the buffer where
-    * they live will not be created and cannot make it resident */
-   if (desc->sets_not_resident)
-      kk_make_descriptor_resources_resident(cmd,
-                                            VK_PIPELINE_BIND_POINT_GRAPHICS);
    if (desc->root_dirty)
       kk_upload_descriptor_root(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS);
 
diff --git a/src/kosmickrisp/vulkan/kk_cmd_meta.c b/src/kosmickrisp/vulkan/kk_cmd_meta.c
index debf9cd467c..c889fe19a0d 100644
--- a/src/kosmickrisp/vulkan/kk_cmd_meta.c
+++ b/src/kosmickrisp/vulkan/kk_cmd_meta.c
@@ -35,8 +35,7 @@ kk_cmd_bind_map_buffer(struct vk_command_buffer *vk_cmd,
    buffer->mtl_handle = bo->map;
    buffer->vk.device_address = bo->gpu;
    *map_out = bo->cpu;
-   mtl_compute_use_resource(cmd->encoder->main.encoder, buffer->mtl_handle,
-                            MTL_RESOURCE_USAGE_WRITE | MTL_RESOURCE_USAGE_READ);
+
    return VK_SUCCESS;
 }
 
@@ -147,11 +146,9 @@ kk_meta_end(struct kk_cmd_buffer *cmd, struct kk_meta_save *save,
       desc->sets[0] = save->desc0;
       desc->root.sets[0] = save->desc0->addr;
       desc->set_sizes[0] = save->desc0->size;
-      desc->sets_not_resident |= BITFIELD_BIT(0);
       desc->push_dirty &= ~BITFIELD_BIT(0);
    } else if (save->has_push_desc0) {
       desc->push[0] = save->push_desc0;
-      desc->sets_not_resident |= BITFIELD_BIT(0);
       desc->push_dirty |= BITFIELD_BIT(0);
    }
 
@@ -197,13 +194,10 @@ kk_CmdFillBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer,
                  VkDeviceSize dstOffset, VkDeviceSize dstRange, uint32_t data)
 {
    VK_FROM_HANDLE(kk_cmd_buffer, cmd, commandBuffer);
-   VK_FROM_HANDLE(kk_buffer, buf, dstBuffer);
    struct kk_device *dev = kk_cmd_buffer_device(cmd);
 
    struct kk_meta_save save;
    kk_meta_begin(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
-   mtl_compute_use_resource(kk_compute_encoder(cmd), buf->mtl_handle,
-                            MTL_RESOURCE_USAGE_WRITE);
    vk_meta_fill_buffer(&cmd->vk, &dev->meta, dstBuffer, dstOffset, dstRange,
                        data);
    kk_meta_end(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
@@ -215,13 +209,10 @@ kk_CmdUpdateBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer,
                    const void *pData)
 {
    VK_FROM_HANDLE(kk_cmd_buffer, cmd, commandBuffer);
-   VK_FROM_HANDLE(kk_buffer, buf, dstBuffer);
    struct kk_device *dev = kk_cmd_buffer_device(cmd);
 
    struct kk_meta_save save;
    kk_meta_begin(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
-   mtl_compute_use_resource(kk_compute_encoder(cmd), buf->mtl_handle,
-                            MTL_RESOURCE_USAGE_WRITE);
    vk_meta_update_buffer(&cmd->vk, &dev->meta, dstBuffer, dstOffset, dstRange,
                          pData);
    kk_meta_end(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
diff --git a/src/kosmickrisp/vulkan/kk_descriptor_set.c b/src/kosmickrisp/vulkan/kk_descriptor_set.c
index 0637fddc812..090c1b055f4 100644
--- a/src/kosmickrisp/vulkan/kk_descriptor_set.c
+++ b/src/kosmickrisp/vulkan/kk_descriptor_set.c
@@ -579,7 +579,6 @@ kk_descriptor_set_create(struct kk_device *dev, struct kk_descriptor_pool *pool,
          vk_object_free(&dev->vk, NULL, set);
          return result;
       }
-      set->mtl_descriptor_buffer = pool->bo->map;
    }
 
    vk_descriptor_set_layout_ref(&layout->vk);
diff --git a/src/kosmickrisp/vulkan/kk_descriptor_set.h b/src/kosmickrisp/vulkan/kk_descriptor_set.h
index d1652448a41..741fcc73b5a 100644
--- a/src/kosmickrisp/vulkan/kk_descriptor_set.h
+++ b/src/kosmickrisp/vulkan/kk_descriptor_set.h
@@ -41,7 +41,6 @@ struct kk_descriptor_set {
    struct list_head link;
 
    struct kk_descriptor_set_layout *layout;
-   mtl_resource *mtl_descriptor_buffer;
    void *mapped_ptr;
    uint64_t addr;
    uint32_t size;
@@ -64,9 +63,6 @@ kk_descriptor_set_addr(const struct kk_descriptor_set *set)
 struct kk_push_descriptor_set {
    uint8_t data[KK_PUSH_DESCRIPTOR_SET_SIZE];
    struct kk_descriptor_set_layout *layout;
-   mtl_resource *mtl_descriptor_buffer;
-   uint32_t resource_count;
-   mtl_resource *mtl_resources[];
 };
 
 void kk_push_descriptor_set_update(struct kk_push_descriptor_set *push_set,
diff --git a/src/kosmickrisp/vulkan/kk_descriptor_set_layout.c b/src/kosmickrisp/vulkan/kk_descriptor_set_layout.c
index 21004dc818d..97c76a4d831 100644
--- a/src/kosmickrisp/vulkan/kk_descriptor_set_layout.c
+++ b/src/kosmickrisp/vulkan/kk_descriptor_set_layout.c
@@ -176,7 +176,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
    uint32_t buffer_size = 0;
    uint32_t max_variable_descriptor_size = 0;
    uint8_t dynamic_buffer_count = 0;
-   uint32_t total_descriptor_count = 0u;
    for (uint32_t b = 0; b < num_bindings; b++) {
       /* We stashed the pCreateInfo->pBindings[] index (plus one) in the
        * immutable_samplers pointer.  Check for NULL (empty binding) and then
@@ -195,7 +194,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
          continue;
 
       layout->binding[b].type = binding->descriptorType;
-      layout->binding[b].mtl_resources_index = total_descriptor_count;
       layout->descriptor_count += binding->descriptorCount;
 
       if (binding_flags_info && binding_flags_info->bindingCount > 0) {
@@ -250,7 +248,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
 
       stride *= max_plane_count;
       layout->binding[b].count_per_element = max_plane_count;
-      total_descriptor_count += max_plane_count * binding->descriptorCount;
 
       if (stride > 0) {
          assert(stride <= UINT8_MAX);
diff --git a/src/kosmickrisp/vulkan/kk_descriptor_set_layout.h b/src/kosmickrisp/vulkan/kk_descriptor_set_layout.h
index 14ecda0f62a..a993a4a2dd2 100644
--- a/src/kosmickrisp/vulkan/kk_descriptor_set_layout.h
+++ b/src/kosmickrisp/vulkan/kk_descriptor_set_layout.h
@@ -38,9 +38,6 @@ struct kk_descriptor_set_binding_layout {
    /* Offset into the descriptor buffer where this descriptor lives */
    uint32_t offset;
 
-   /* Offset to the mtl_resource_ids array where this descriptor stores them */
-   uint32_t mtl_resources_index;
-
    /* Stride between array elements in the descriptor buffer */
    uint8_t stride;
 
diff --git a/src/kosmickrisp/vulkan/kk_device.c b/src/kosmickrisp/vulkan/kk_device.c
index bb05a30cbd5..93e959ce746 100644
--- a/src/kosmickrisp/vulkan/kk_device.c
+++ b/src/kosmickrisp/vulkan/kk_device.c
@@ -205,10 +205,10 @@ kk_CreateDevice(VkPhysicalDevice physicalDevice,
    dev->vk.command_buffer_ops = &kk_cmd_buffer_ops;
    dev->vk.command_dispatch_table = &dev->vk.dispatch_table;
 
-   /* Buffer to use as null descriptor */
-   result = kk_alloc_bo(dev, &dev->vk.base, sizeof(uint64_t) * 8, 8u,
-                        &dev->null_descriptor);
-   if (result != VK_SUCCESS)
+   /* We need to initialize the device residency set before any bo is created. */
+   simple_mtx_init(&dev->residency_set.mutex, mtx_plain);
+   dev->residency_set.handle = mtl_new_residency_set(dev->mtl_handle);
+   if (dev->residency_set.handle == NULL)
       goto fail_init;
 
    result =
@@ -233,10 +233,6 @@ kk_CreateDevice(VkPhysicalDevice physicalDevice,
    if (result != VK_SUCCESS)
       goto fail_sampler_heap;
 
-   simple_mtx_init(&dev->user_residency_set.mutex, mtx_plain);
-   dev->user_residency_set.residency_set =
-      mtl_new_residency_set(dev->mtl_handle);
-
    kk_parse_device_environment_options(dev);
 
    *pDevice = kk_device_to_handle(dev);
@@ -252,7 +248,8 @@ fail_meta:
 fail_mem_cache:
    kk_queue_finish(dev, &dev->queue);
 fail_vab_memory:
-   kk_destroy_bo(dev, dev->null_descriptor);
+   mtl_release(dev->residency_set.handle);
+   simple_mtx_destroy(&dev->residency_set.mutex);
 fail_init:
    vk_device_finish(&dev->vk);
 fail_alloc:
@@ -268,25 +265,24 @@ kk_DestroyDevice(VkDevice _device, const VkAllocationCallbacks *pAllocator)
    if (!dev)
       return;
 
+   /* End capture before we start releasing resources. Otherwise, Metal capture
+    * may run into issues. */
+   if (dev->gpu_capture_enabled) {
+      mtl_stop_gpu_capture();
+   }
+
    /* Meta first since it may destroy Vulkan objects */
    kk_device_finish_meta(dev);
-
-   /* Need to end the residency otherwise stopping a capture crashes the
-    * program... */
-   mtl_residency_set_end_residency(dev->user_residency_set.residency_set);
-   mtl_release(dev->user_residency_set.residency_set);
-   simple_mtx_destroy(&dev->user_residency_set.mutex);
    kk_device_finish_lib(dev);
    kk_query_table_finish(dev, &dev->occlusion_queries);
    kk_destroy_sampler_heap(dev, &dev->samplers);
 
-   kk_queue_finish(dev, &dev->queue);
-   kk_destroy_bo(dev, dev->null_descriptor);
-   vk_device_finish(&dev->vk);
+   /* Release the residency set last once all BOs are released. */
+   mtl_release(dev->residency_set.handle);
+   simple_mtx_destroy(&dev->residency_set.mutex);
 
-   if (dev->gpu_capture_enabled) {
-      mtl_stop_gpu_capture();
-   }
+   kk_queue_finish(dev, &dev->queue);
+   vk_device_finish(&dev->vk);
 
    vk_free(&dev->vk.alloc, dev);
 }
@@ -358,28 +354,26 @@ kk_GetDeviceProcAddr(VkDevice _device, const char *pName)
 }
 
 void
-kk_device_add_user_heap(struct kk_device *dev, mtl_heap *heap)
+kk_device_add_heap_to_residency_set(struct kk_device *dev, mtl_heap *heap)
 {
-   simple_mtx_lock(&dev->user_residency_set.mutex);
-   mtl_residency_set_add_allocation(dev->user_residency_set.residency_set,
-                                    heap);
-   simple_mtx_unlock(&dev->user_residency_set.mutex);
+   simple_mtx_lock(&dev->residency_set.mutex);
+   mtl_residency_set_add_allocation(dev->residency_set.handle, heap);
+   simple_mtx_unlock(&dev->residency_set.mutex);
 }
 
 void
-kk_device_remove_user_heap(struct kk_device *dev, mtl_heap *heap)
+kk_device_remove_heap_from_residency_set(struct kk_device *dev, mtl_heap *heap)
 {
-   simple_mtx_lock(&dev->user_residency_set.mutex);
-   mtl_residency_set_remove_allocation(dev->user_residency_set.residency_set,
-                                       heap);
-   simple_mtx_unlock(&dev->user_residency_set.mutex);
+   simple_mtx_lock(&dev->residency_set.mutex);
+   mtl_residency_set_remove_allocation(dev->residency_set.handle, heap);
+   simple_mtx_unlock(&dev->residency_set.mutex);
 }
 
 void
 kk_device_make_resources_resident(struct kk_device *dev)
 {
-   simple_mtx_lock(&dev->user_residency_set.mutex);
-   mtl_residency_set_commit(dev->user_residency_set.residency_set);
-   mtl_residency_set_request_residency(dev->user_residency_set.residency_set);
-   simple_mtx_unlock(&dev->user_residency_set.mutex);
+   simple_mtx_lock(&dev->residency_set.mutex);
+   mtl_residency_set_commit(dev->residency_set.handle);
+   mtl_residency_set_request_residency(dev->residency_set.handle);
+   simple_mtx_unlock(&dev->residency_set.mutex);
 }
diff --git a/src/kosmickrisp/vulkan/kk_device.h b/src/kosmickrisp/vulkan/kk_device.h
index d853fdc2a2f..77e84eb0b20 100644
--- a/src/kosmickrisp/vulkan/kk_device.h
+++ b/src/kosmickrisp/vulkan/kk_device.h
@@ -32,9 +32,9 @@ enum kk_device_lib_pipeline {
    KK_LIB_COUNT,
 };
 
-struct kk_user_residency_set {
+struct kk_residency_set {
    simple_mtx_t mutex;
-   mtl_residency_set *residency_set;
+   mtl_residency_set *handle;
 };
 
 struct mtl_sampler_packed {
@@ -84,14 +84,12 @@ struct kk_device {
     * commands due to Metal limitations */
    struct vk_device_dispatch_table exposed_dispatch_table;
 
-   struct kk_bo *null_descriptor;
-
    struct kk_sampler_heap samplers;
    struct kk_query_table occlusion_queries;
 
    /* Track all heaps the user allocated so we can set them all as resident when
     * recording as required by Metal. */
-   struct kk_user_residency_set user_residency_set;
+   struct kk_residency_set residency_set;
 
    mtl_compute_pipeline_state *lib_pipelines[KK_LIB_COUNT];
 
@@ -123,8 +121,9 @@ VkResult kk_device_init_meta(struct kk_device *dev);
 void kk_device_finish_meta(struct kk_device *dev);
 VkResult kk_device_init_lib(struct kk_device *dev);
 void kk_device_finish_lib(struct kk_device *dev);
-void kk_device_add_user_heap(struct kk_device *dev, mtl_heap *heap);
-void kk_device_remove_user_heap(struct kk_device *dev, mtl_heap *heap);
+void kk_device_add_heap_to_residency_set(struct kk_device *dev, mtl_heap *heap);
+void kk_device_remove_heap_from_residency_set(struct kk_device *dev,
+                                              mtl_heap *heap);
 void kk_device_make_resources_resident(struct kk_device *dev);
 
 /* Required to create a sampler */
diff --git a/src/kosmickrisp/vulkan/kk_device_memory.c b/src/kosmickrisp/vulkan/kk_device_memory.c
index e020aa979f4..13e1dcc6595 100644
--- a/src/kosmickrisp/vulkan/kk_device_memory.c
+++ b/src/kosmickrisp/vulkan/kk_device_memory.c
@@ -94,6 +94,7 @@ kk_AllocateMemory(VkDevice device, const VkMemoryAllocateInfo *pAllocateInfo,
       mem->bo->gpu = mtl_buffer_get_gpu_address(mem->bo->map);
       mem->bo->cpu = mtl_get_contents(mem->bo->map);
       mem->bo->size_B = mtl_heap_get_size(mem->bo->mtl_handle);
+      kk_device_add_heap_to_residency_set(dev, mem->bo->mtl_handle);
    } else {
       result =
          kk_alloc_bo(dev, &dev->vk.base, aligned_size, alignment, &mem->bo);
@@ -104,8 +105,6 @@ kk_AllocateMemory(VkDevice device, const VkMemoryAllocateInfo *pAllocateInfo,
    struct kk_memory_heap *heap = &pdev->mem_heaps[type->heapIndex];
    p_atomic_add(&heap->used, mem->bo->size_B);
 
-   kk_device_add_user_heap(dev, mem->bo->mtl_handle);
-
    *pMem = kk_device_memory_to_handle(mem);
 
    return VK_SUCCESS;
@@ -126,8 +125,6 @@ kk_FreeMemory(VkDevice device, VkDeviceMemory _mem,
    if (!mem)
       return;
 
-   kk_device_remove_user_heap(dev, mem->bo->mtl_handle);
-
    const VkMemoryType *type = &pdev->mem_types[mem->vk.memory_type_index];
    struct kk_memory_heap *heap = &pdev->mem_heaps[type->heapIndex];
    p_atomic_add(&heap->used, -((int64_t)mem->bo->size_B));
diff --git a/src/kosmickrisp/vulkan/kk_encoder.c b/src/kosmickrisp/vulkan/kk_encoder.c
index 76eb5ad5806..e65aedc5443 100644
--- a/src/kosmickrisp/vulkan/kk_encoder.c
+++ b/src/kosmickrisp/vulkan/kk_encoder.c
@@ -39,7 +39,6 @@ kk_encoder_init(mtl_device *device, struct kk_queue *queue,
    kk_encoder_start_internal(&enc->pre_gfx, device, queue->pre_gfx.mtl_handle);
    enc->event = mtl_new_event(device);
    enc->imm_writes = UTIL_DYNARRAY_INIT;
-   enc->resident_buffers = UTIL_DYNARRAY_INIT;
    enc->copy_query_pool_result_infos = UTIL_DYNARRAY_INIT;
 
    *encoder = enc;
@@ -180,12 +179,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
       if (!bo)
          return;
       memcpy(bo->cpu, enc->imm_writes.data, enc->imm_writes.size);
-      uint32_t buffer_count =
-         util_dynarray_num_elements(&enc->resident_buffers, mtl_buffer *);
-      mtl_compute_use_resource(compute, bo->map, MTL_RESOURCE_USAGE_READ);
-      mtl_compute_use_resources(
-         compute, enc->resident_buffers.data, buffer_count,
-         MTL_RESOURCE_USAGE_READ | MTL_RESOURCE_USAGE_WRITE);
       struct kk_imm_write_push push_data = {
          .buffer_address = bo->gpu,
          .count = count,
@@ -193,7 +186,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
       kk_cmd_dispatch_pipeline(cmd, compute,
                                kk_device_lib_pipeline(dev, KK_LIB_IMM_WRITE),
                                &push_data, sizeof(push_data), count, 1, 1);
-      enc->resident_buffers.size = 0u;
       enc->imm_writes.size = 0u;
    }
 
@@ -201,11 +193,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
                                       struct kk_copy_query_pool_results_info);
    if (count != 0u) {
       mtl_compute_encoder *compute = kk_compute_encoder(cmd);
-      uint32_t buffer_count =
-         util_dynarray_num_elements(&enc->resident_buffers, mtl_buffer *);
-      mtl_compute_use_resources(
-         compute, enc->resident_buffers.data, buffer_count,
-         MTL_RESOURCE_USAGE_READ | MTL_RESOURCE_USAGE_WRITE);
 
       for (uint32_t i = 0u; i < count; ++i) {
          struct kk_copy_query_pool_results_info *push_data =
@@ -216,7 +203,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
             cmd, compute, kk_device_lib_pipeline(dev, KK_LIB_COPY_QUERY),
             push_data, sizeof(*push_data), push_data->query_count, 1, 1);
       }
-      enc->resident_buffers.size = 0u;
       enc->copy_query_pool_result_infos.size = 0u;
    }
 
@@ -292,7 +278,6 @@ kk_post_execution_release(void *data)
    kk_post_execution_release_internal(&encoder->pre_gfx);
    mtl_release(encoder->event);
    util_dynarray_fini(&encoder->imm_writes);
-   util_dynarray_fini(&encoder->resident_buffers);
    util_dynarray_fini(&encoder->copy_query_pool_result_infos);
    free(encoder);
 }
@@ -416,11 +401,6 @@ kk_encoder_render_triangle_fan_common(struct kk_cmd_buffer *cmd,
    info->out_el_size_B = out_el_size_B;
    info->flatshade_first = true;
    mtl_compute_encoder *encoder = kk_encoder_pre_gfx_encoder(cmd);
-   if (index)
-      mtl_compute_use_resource(encoder, index, MTL_RESOURCE_USAGE_READ);
-   mtl_compute_use_resource(encoder, indirect, MTL_RESOURCE_USAGE_READ);
-   mtl_compute_use_resource(encoder, index_buffer->map,
-                            MTL_RESOURCE_USAGE_WRITE);
 
    struct kk_device *dev = kk_cmd_buffer_device(cmd);
    kk_cmd_dispatch_pipeline(cmd, encoder,
diff --git a/src/kosmickrisp/vulkan/kk_encoder.h b/src/kosmickrisp/vulkan/kk_encoder.h
index 5c45b87ecdd..d216bdbf592 100644
--- a/src/kosmickrisp/vulkan/kk_encoder.h
+++ b/src/kosmickrisp/vulkan/kk_encoder.h
@@ -70,9 +70,6 @@ struct kk_encoder {
    /* uint64_t pairs with first being the address, second being the value to
     * write */
    struct util_dynarray imm_writes;
-   /* mtl_buffers (destination buffers) so we can make them resident before the
-    * dispatch */
-   struct util_dynarray resident_buffers;
    /* Array of kk_copy_quer_pool_results_info structs */
    struct util_dynarray copy_query_pool_result_infos;
 };
diff --git a/src/kosmickrisp/vulkan/kk_query_pool.c b/src/kosmickrisp/vulkan/kk_query_pool.c
index eefa33f2deb..0d58cfb4a13 100644
--- a/src/kosmickrisp/vulkan/kk_query_pool.c
+++ b/src/kosmickrisp/vulkan/kk_query_pool.c
@@ -418,10 +418,6 @@ kk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool,
    };
 
    util_dynarray_append(&cmd->encoder->copy_query_pool_result_infos, info);
-   util_dynarray_append(&cmd->encoder->resident_buffers, dst_buf->mtl_handle);
-   util_dynarray_append(&cmd->encoder->resident_buffers, pool->bo->map);
-   util_dynarray_append(&cmd->encoder->resident_buffers,
-                        dev->occlusion_queries.bo->map);
    /* If we are not mid encoder, just upload the writes */
    if (cmd->encoder->main.last_used == KK_ENC_NONE)
       upload_queue_writes(cmd);
diff --git a/src/kosmickrisp/vulkan/kk_queue.c b/src/kosmickrisp/vulkan/kk_queue.c
index 1e2ea723cc9..2a8b2b22237 100644
--- a/src/kosmickrisp/vulkan/kk_queue.c
+++ b/src/kosmickrisp/vulkan/kk_queue.c
@@ -31,11 +31,6 @@ kk_queue_submit(struct vk_queue *vk_queue, struct vk_queue_submit *submit)
    if (result != VK_SUCCESS)
       return result;
 
-   /* Ensure any changes to residency are propagated before we submit any work.
-    * All resources should have been allocated before submission. Otherwise,
-    * users are playing with fire. */
-   kk_device_make_resources_resident(dev);
-
    /* Chain with previous sumbission */
    if (queue->wait_fence) {
       util_dynarray_append(&encoder->main.fences, queue->wait_fence);
@@ -78,6 +73,11 @@ kk_queue_submit(struct vk_queue *vk_queue, struct vk_queue_submit *submit)
                               signal->signal_value);
    }
 
+   /* Ensure any changes to residency are propagated before we submit any work.
+    * All resources should have been allocated before submission. Otherwise,
+    * users are playing with fire. */
+   kk_device_make_resources_resident(dev);
+
    /* Steal the last fence to chain with the next submission */
    if (util_dynarray_num_elements(&encoder->main.fences, mtl_fence *) > 0)
       queue->wait_fence = util_dynarray_pop(&encoder->main.fences, mtl_fence *);