Cycles: Metal: Integrate MTLResidencySets for explicit memory management

Introduce `MTLResidencySet` to explicitly manage GPU memory residency on macOS 15.0+ devices.

This provides the Metal/GPU driver with a clear list of resources to optimise memory handling, potentially improving performance by reducing overhead from hundreds of `useResource` calls.

Memory allocation and deallocation operations are now routed through new wrapper functions that conditionally add or remove resources from the residency set. Explicit `useResource` calls are bypassed when residency sets are active, as resource residency is now managed through the `MTLResidencySet` API. A debug flag is added to enable or disable this feature at runtime (via env var `CYCLES_METAL_RESIDENCY_SETS=0`).

Pull Request: https://projects.blender.org/blender/blender/pulls/158558

intern/cycles/device/metal/bvh.mm

@@ -161,14 +161,12 @@ void BVHMetal::set_accel_struct(id<MTLAccelerationStructure> new_accel_struct)
 {
   if (@available(macos 12.0, *)) {
     if (accel_struct) {
-      device->stats.mem_free(accel_struct.allocatedSize);
       [accel_struct release];
       accel_struct = nil;
     }
 
     if (new_accel_struct) {
       accel_struct = new_accel_struct;
-      device->stats.mem_alloc(accel_struct.allocatedSize);
     }
   }
 }

intern/cycles/device/metal/device_impl.h

@@ -53,7 +53,20 @@ class MetalDevice : public Device {
 
   API_AVAILABLE(macos(11.0))
   id<MTLAccelerationStructure> accel_struct = nil;
-  /* --------------------------------------------------- */
+
+  /* Residency sets -----------------------------------*/
+  void prepare_residency();
+  void metal_mem_alloc(id<MTLResource> allocation);
+  void metal_mem_free(id<MTLResource> allocation);
+
+  bool mtlResidencySet_enabled = false;
+#  if defined(MAC_OS_VERSION_15_0)
+  API_AVAILABLE(macos(15.0), ios(18.0))
+  id<MTLResidencySet> mtlResidencySet = nil;
+  bool mtlResidencySet_dirty = false;
+  /* Guards mtlResidencySet mutations (may be reached from multiple threads). */
+  std::mutex mtlResidencySet_mutex;
+#  endif
 
   uint kernel_features = 0;
   bool using_nanovdb = false;

intern/cycles/device/metal/device_impl.mm

@@ -163,16 +163,6 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
                  kernel_type_as_string(
                      (MetalPipelineType)min((int)kernel_specialization_level, (int)PSO_NUM - 1)));
 
-    image_bindings = [mtlDevice newBufferWithLength:8192 options:MTLResourceStorageModeShared];
-    stats.mem_alloc(image_bindings.allocatedSize);
-
-    launch_params_buffer = [mtlDevice newBufferWithLength:sizeof(KernelParamsMetal)
-                                                  options:MTLResourceStorageModeShared];
-    stats.mem_alloc(sizeof(KernelParamsMetal));
-
-    /* Cache unified pointer so we can write kernel params directly in place. */
-    launch_params = (KernelParamsMetal *)launch_params_buffer.contents;
-
     /* Command queue for path-tracing work on the GPU. In a situation where multiple
      * MetalDeviceQueues are spawned from one MetalDevice, they share the same MTLCommandQueue.
      * This is thread safe and just as performant as each having their own instance. It also
@@ -181,6 +171,43 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
 
     /* Command queue for non-tracing work on the GPU. */
     mtlGeneralCommandQueue = [mtlDevice newCommandQueue];
+
+#  if defined(MAC_OS_VERSION_15_0)
+    if (@available(macos 15.0, *)) {
+      if (DebugFlags().metal.use_residency_sets_if_available) {
+        /* Use a residency set to declare all rendering resources up front, avoiding
+         * the overhead of per-encoder useResource calls on every dispatch. */
+        MTLResidencySetDescriptor *residency_set_desc = [[MTLResidencySetDescriptor alloc] init];
+        residency_set_desc.label = @"CyclesResidencySet";
+        residency_set_desc.initialCapacity = 512;
+        NSError *error = nil;
+        mtlResidencySet = [mtlDevice newResidencySetWithDescriptor:residency_set_desc
+                                                             error:&error];
+        [residency_set_desc release];
+
+        /* Only enable residency sets if creation succeeded. Otherwise we fall back to the
+         * per-encoder useResource path. */
+        if (mtlResidencySet) {
+          mtlResidencySet_enabled = true;
+          [mtlComputeCommandQueue addResidencySet:mtlResidencySet];
+        }
+        else {
+          metal_printf("Failed to create residency set: %s",
+                       [[error localizedDescription] UTF8String]);
+        }
+      }
+    }
+#  endif
+
+    image_bindings = [mtlDevice newBufferWithLength:8192 options:MTLResourceStorageModeShared];
+    metal_mem_alloc(image_bindings);
+
+    launch_params_buffer = [mtlDevice newBufferWithLength:sizeof(KernelParamsMetal)
+                                                  options:MTLResourceStorageModeShared];
+    metal_mem_alloc(launch_params_buffer);
+
+    /* Cache unified pointer so we can write kernel params directly in place. */
+    launch_params = (KernelParamsMetal *)launch_params_buffer.contents;
   }
 }
 
@@ -195,18 +222,25 @@ MetalDevice::~MetalDevice()
 
   /* Release textures that weren't already freed by tex_free. */
   for (int res = 0; res < image_info.size(); res++) {
-    [image_info_id_map[res] release];
+    metal_mem_free(image_info_id_map[res]);
     image_info_id_map[res] = nil;
   }
 
+  /* Queue resources for release, then run flush_delayed_free_list(). */
   free_bvh();
+  metal_mem_free(launch_params_buffer);
+  metal_mem_free(image_bindings);
+  image_info.free();
   flush_delayed_free_list();
 
-  stats.mem_free(sizeof(KernelParamsMetal));
-  [launch_params_buffer release];
-
-  stats.mem_free(image_bindings.allocatedSize);
-  [image_bindings release];
+#  if defined(MAC_OS_VERSION_15_0)
+  if (@available(macos 15.0, *)) {
+    if (mtlResidencySet) {
+      [mtlResidencySet endResidency];
+      [mtlResidencySet release];
+    }
+  }
+#  endif
 
   [mtlComputeCommandQueue release];
   [mtlGeneralCommandQueue release];
@@ -214,8 +248,62 @@ MetalDevice::~MetalDevice()
     [mtlCounterSampleBuffer release];
   }
   [mtlDevice release];
+}
 
-  image_info.free();
+void MetalDevice::metal_mem_alloc(id<MTLResource> allocation)
+{
+  if (allocation) {
+    stats.mem_alloc(allocation.allocatedSize);
+#  if defined(MAC_OS_VERSION_15_0)
+    if (@available(macos 15.0, *)) {
+      if (mtlResidencySet) {
+        std::lock_guard<std::mutex> residency_lock(mtlResidencySet_mutex);
+        [mtlResidencySet addAllocation:allocation];
+        mtlResidencySet_dirty = true;
+      }
+    }
+#  endif
+  }
+}
+
+void MetalDevice::metal_mem_free(id<MTLResource> allocation)
+{
+  if (allocation) {
+    stats.mem_free(allocation.allocatedSize);
+
+    std::lock_guard<std::recursive_mutex> lock(metal_mem_map_mutex);
+#  if defined(MAC_OS_VERSION_15_0)
+    /* Remove from the residency set immediately, but don't commit until next enqueue. A resource
+     * can be repurposed (e.g. a BVH refit) so a deferred removal can spuriously swap the
+     * remove-then-add to be an add-then-remove. */
+    if (@available(macos 15.0, *)) {
+      if (mtlResidencySet) {
+        std::lock_guard<std::mutex> residency_lock(mtlResidencySet_mutex);
+        [mtlResidencySet removeAllocation:allocation];
+        mtlResidencySet_dirty = true;
+      }
+    }
+#  endif
+
+    /* Defer the actual [release] until flush_delayed_free_list(), so the object stays alive for
+     * any command buffer still referencing it. */
+    delayed_free_list.push_back(allocation);
+  }
+}
+
+void MetalDevice::prepare_residency()
+{
+#  if defined(MAC_OS_VERSION_15_0)
+  if (@available(macos 15.0, *)) {
+    if (mtlResidencySet) {
+      std::lock_guard<std::mutex> residency_lock(mtlResidencySet_mutex);
+      if (mtlResidencySet_dirty) {
+        mtlResidencySet_dirty = false;
+        [mtlResidencySet commit];
+      }
+    }
+  }
+#  endif
 }
 
 bool MetalDevice::support_device(const uint /*kernel_features*/)
@@ -559,7 +647,6 @@ bool MetalDevice::is_texture(const KernelImageInfo &info)
 
 void MetalDevice::erase_allocation(device_memory &mem)
 {
-  stats.mem_free(mem.device_size);
   mem.device_pointer = 0;
   mem.device_size = 0;
 
@@ -612,7 +699,7 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
               << string_human_readable_size(mem.memory_size()) << ")";
 
     mem.device_size = metal_buffer.allocatedSize;
-    stats.mem_alloc(mem.device_size);
+    metal_mem_alloc(metal_buffer);
 
     metal_buffer.label = [NSString stringWithFormat:@"%s", mem.log_name().c_str()];
 
@@ -702,7 +789,7 @@ void MetalDevice::generic_free(device_memory &mem)
     mem.shared_pointer = nullptr;
 
     /* Free device memory. */
-    delayed_free_list.push_back(mmem.mtlBuffer);
+    metal_mem_free(mmem.mtlBuffer);
     mmem.mtlBuffer = nil;
   }
 
@@ -1124,7 +1211,7 @@ void MetalDevice::image_alloc(device_image &mem)
 
     mem.device_pointer = (device_ptr)mtlTexture;
     mem.device_size = size;
-    stats.mem_alloc(size);
+    metal_mem_alloc(mtlTexture);
 
     std::lock_guard<std::recursive_mutex> lock(metal_mem_map_mutex);
     unique_ptr<MetalMem> mmem = make_unique<MetalMem>();
@@ -1143,13 +1230,12 @@ void MetalDevice::image_alloc(device_image &mem)
       ssize_t min_buffer_length = sizeof(void *) * image_info.size();
       if (!image_bindings || (image_bindings.length < min_buffer_length)) {
         if (image_bindings) {
-          delayed_free_list.push_back(image_bindings);
-          stats.mem_free(image_bindings.allocatedSize);
+          metal_mem_free(image_bindings);
         }
         image_bindings = [mtlDevice newBufferWithLength:min_buffer_length
                                                 options:MTLResourceStorageModeShared];
 
-        stats.mem_alloc(image_bindings.allocatedSize);
+        metal_mem_alloc(image_bindings);
       }
     }
 
@@ -1197,7 +1283,7 @@ void MetalDevice::image_free(device_image &mem)
     MetalMem &mmem = *metal_mem_map.at(&mem);
 
     /* Free bindless texture. */
-    delayed_free_list.push_back(mmem.mtlTexture);
+    metal_mem_free(mmem.mtlTexture);
     mmem.mtlTexture = nil;
     erase_allocation(mem);
   }
@@ -1265,19 +1351,21 @@ void MetalDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
 
 void MetalDevice::free_bvh()
 {
+  /* metal_mem_free defers the actual release via delayed_free_list,
+   * since the old BVH may still be referenced by an in-flight command buffer. */
   for (id<MTLAccelerationStructure> &blas : unique_blas_array) {
-    [blas release];
+    metal_mem_free(blas);
   }
   unique_blas_array.clear();
   blas_array.clear();
 
   if (blas_buffer) {
-    [blas_buffer release];
+    metal_mem_free(blas_buffer);
     blas_buffer = nil;
   }
 
   if (accel_struct) {
-    [accel_struct release];
+    metal_mem_free(accel_struct);
     accel_struct = nil;
   }
 }
@@ -1294,15 +1382,20 @@ void MetalDevice::update_bvh(BVHMetal *bvh_metal)
   unique_blas_array = bvh_metal->unique_blas_array;
   blas_array = bvh_metal->blas_array;
 
+  /* Memory tracking and residency are managed here (not in BVHMetal::set_accel_struct)
+   * to pair with free_bvh and reflect actual device ownership. */
+  metal_mem_alloc(accel_struct);
+
   [accel_struct retain];
   for (id<MTLAccelerationStructure> &blas : unique_blas_array) {
     [blas retain];
+    metal_mem_alloc(blas);
   }
 
   // Allocate required buffers for BLAS array.
   uint64_t buffer_size = blas_array.size() * sizeof(uint64_t);
   blas_buffer = [mtlDevice newBufferWithLength:buffer_size options:MTLResourceStorageModeShared];
-  stats.mem_alloc(blas_buffer.allocatedSize);
+  metal_mem_alloc(blas_buffer);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/metal/kernel.h

@@ -104,6 +104,7 @@ class MetalDispatchPipeline {
 
   int pipeline_id = -1;
 
+  MetalDevice *metal_device = nullptr;
   MetalPipelineType pso_type;
   id<MTLComputePipelineState> pipeline = nil;
   int num_threads_per_block = 0;

intern/cycles/device/metal/kernel.mm

@@ -473,7 +473,8 @@ void MetalDispatchPipeline::free_intersection_function_tables()
 {
   for (int table = 0; table < METALRT_TABLE_NUM; table++) {
     if (intersection_func_table[table]) {
-      [intersection_func_table[table] release];
+      /* Add the table to the delayed free list of the device that created it. */
+      metal_device->metal_mem_free(intersection_func_table[table]);
       intersection_func_table[table] = nil;
     }
   }
@@ -486,6 +487,7 @@ MetalDispatchPipeline::~MetalDispatchPipeline()
 
 bool MetalDispatchPipeline::update(MetalDevice *metal_device, DeviceKernel kernel)
 {
+  this->metal_device = metal_device;
   const MetalKernelPipeline *best_pipeline = MetalDeviceKernels::get_best_pipeline(metal_device,
                                                                                    kernel);
   if (!best_pipeline) {
@@ -520,6 +522,15 @@ bool MetalDispatchPipeline::update(MetalDevice *metal_device, DeviceKernel kerne
               functionHandleWithFunction:best_pipeline->table_functions[table][i]];
           [intersection_func_table[table] setFunction:handle atIndex:i];
         }
+
+        /* Bind launch_params into the intersection function table once, when the table is
+         * (re)created. launch_params_buffer is allocated once and never moves, and the binding
+         * persists on the table, so there's no need to rebind it on every dispatch. */
+        [intersection_func_table[table] setBuffer:metal_device->launch_params_buffer
+                                           offset:0
+                                          atIndex:1];
+
+        metal_device->metal_mem_alloc(intersection_func_table[table]);
       }
     }
   }

intern/cycles/device/metal/queue.h

@@ -56,7 +56,7 @@ class MetalDeviceQueue : public DeviceQueue {
   void update_capture(DeviceKernel kernel);
   void begin_capture();
   void end_capture();
-  void prepare_resources(DeviceKernel kernel);
+  void prepare_resources();
 
   id<MTLComputeCommandEncoder> get_compute_encoder(DeviceKernel kernel);
   id<MTLBlitCommandEncoder> get_blit_encoder();

intern/cycles/device/metal/queue.mm

@@ -479,10 +479,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       dynamic_bytes_written = round_up(dynamic_bytes_written, size_in_bytes);
       memcpy(dynamic_args + dynamic_bytes_written, args.values[i], size_in_bytes);
       if (args.types[i] == DeviceKernelArguments::POINTER) {
-        if (id<MTLBuffer> buffer = patch_resource(dynamic_args + dynamic_bytes_written)) {
-          [mtlComputeCommandEncoder useResource:buffer
-                                          usage:MTLResourceUsageRead | MTLResourceUsageWrite];
-        }
+        patch_resource(dynamic_args + dynamic_bytes_written);
       }
       dynamic_bytes_written += size_in_bytes;
     }
@@ -511,25 +508,15 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       assert(ancillary_index == ANCILLARY_SLOT_COUNT);
     }
 
-    /* Encode ancillaries */
-    if (metal_device_->use_metalrt) {
-      for (int table = 0; table < METALRT_TABLE_NUM; table++) {
-        if (active_pipeline.intersection_func_table[table]) {
-          [active_pipeline.intersection_func_table[table]
-              setBuffer:metal_device_->launch_params_buffer
-                 offset:0
-                atIndex:1];
-          [mtlComputeCommandEncoder useResource:active_pipeline.intersection_func_table[table]
-                                          usage:MTLResourceUsageRead];
-        }
-      }
-    }
-
     [mtlComputeCommandEncoder setBytes:dynamic_args length:dynamic_bytes_written atIndex:0];
     [mtlComputeCommandEncoder setBuffer:metal_device_->launch_params_buffer offset:0 atIndex:1];
     [mtlComputeCommandEncoder setBytes:ancillary_args length:sizeof(ancillary_args) atIndex:2];
 
-    if (metal_device_->use_metalrt && device_kernel_has_intersection(kernel)) {
+    /* Fallback path in case residency sets aren't supported:
+     * Call useResource for MetalRT resources not covered by prepare_resources(). */
+    if (!metal_device_->mtlResidencySet_enabled && metal_device_->use_metalrt &&
+        device_kernel_has_intersection(kernel))
+    {
       if (@available(macos 12.0, *)) {
 
         if (id<MTLAccelerationStructure> accel_struct = metal_device_->accel_struct) {
@@ -544,6 +531,13 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
                                            usage:MTLResourceUsageRead];
         }
       }
+
+      for (int table = 0; table < METALRT_TABLE_NUM; table++) {
+        if (active_pipeline.intersection_func_table[table]) {
+          [mtlComputeCommandEncoder useResource:active_pipeline.intersection_func_table[table]
+                                          usage:MTLResourceUsageRead];
+        }
+      }
     }
 
     [mtlComputeCommandEncoder setComputePipelineState:active_pipeline.pipeline];
@@ -588,6 +582,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     [mtlComputeCommandEncoder dispatchThreads:size_threads_per_dispatch
                         threadsPerThreadgroup:size_threads_per_threadgroup];
 
+    metal_device_->prepare_residency();
+
     [mtlCommandBuffer_ addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
       /* Enhanced command buffer errors */
       string str;
@@ -787,8 +783,13 @@ void *MetalDeviceQueue::copy_from_device_synchronized(device_memory &mem,
   return (d_ptr) ? reinterpret_cast<MetalDevice::MetalMem *>(d_ptr)->hostPtr : nullptr;
 }
 
-void MetalDeviceQueue::prepare_resources(DeviceKernel /*kernel*/)
+void MetalDeviceQueue::prepare_resources()
 {
+  if (metal_device_->mtlResidencySet_enabled) {
+    /* All resources are already resident — skip per-encoder useResource calls. */
+    return;
+  }
+
   std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
 
   /* declare resource usage */
@@ -828,7 +829,7 @@ id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel
                                                         MTLDispatchTypeSerial)
     {
       /* declare usage of MTLBuffers etc */
-      prepare_resources(kernel);
+      prepare_resources();
 
       return mtlComputeEncoder_;
     }
@@ -865,7 +866,7 @@ id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel
   [mtlComputeEncoder_ setLabel:@(device_kernel_as_string(kernel))];
 
   /* declare usage of MTLBuffers etc */
-  prepare_resources(kernel);
+  prepare_resources();
 
   return mtlComputeEncoder_;
 }

intern/cycles/util/debug.cpp

@@ -84,6 +84,10 @@ void DebugFlags::Metal::reset()
   if (const char *str = getenv("CYCLES_METALRT_PCMI")) {
     use_metalrt_pcmi = (atoi(str) != 0);
   }
+
+  if (const char *str = getenv("CYCLES_METAL_RESIDENCY_SETS")) {
+    use_residency_sets_if_available = (atoi(str) != 0);
+  }
 }
 
 DebugFlags::TextureCache::TextureCache()

intern/cycles/util/debug.h

@@ -101,9 +101,11 @@ class DebugFlags {
     /* Whether async PSO creation is enabled or not. */
     bool use_async_pso_creation = true;
 
-    /* Whether to use per-component motion interpolation.
-     */
+    /* Whether to use per-component motion interpolation. */
     bool use_metalrt_pcmi = true;
+
+    /* Whether to use residency sets. */
+    bool use_residency_sets_if_available = true;
   };
 
   /* Descriptor of Texture Cache feature-set to be used. */