Add Chromium-only Blender WebEngine parity work

blender-5.2.0/intern/cycles/device/metal/bvh.h

@@ -0,0 +1,72 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#ifdef WITH_METAL
+
+#  include "bvh/bvh.h"
+#  include "bvh/params.h"
+#  include "device/memory.h"
+
+#  include <Metal/Metal.h>
+
+CCL_NAMESPACE_BEGIN
+
+class BVHMetal : public BVH {
+ public:
+  API_AVAILABLE(macos(11.0))
+  id<MTLAccelerationStructure> accel_struct = nil;
+
+  API_AVAILABLE(macos(11.0))
+  id<MTLAccelerationStructure> null_BLAS = nil;
+
+  API_AVAILABLE(macos(11.0))
+  vector<id<MTLAccelerationStructure>> blas_array;
+
+  API_AVAILABLE(macos(11.0))
+  vector<id<MTLAccelerationStructure>> unique_blas_array;
+
+  Device *device = nullptr;
+
+  bool motion_blur = false;
+
+  /* Per-component Motion Interpolation in macOS 15. */
+  bool use_pcmi = false;
+
+  bool extended_limits = false;
+
+  bool build(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
+
+  BVHMetal(const BVHParams &params,
+           const vector<Geometry *> &geometry,
+           const vector<Object *> &objects,
+           Device *device);
+  ~BVHMetal() override;
+
+  bool build_BLAS(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
+  bool build_BLAS_mesh(Progress &progress,
+                       id<MTLDevice> device,
+                       id<MTLCommandQueue> queue,
+                       Geometry *const geom,
+                       bool refit);
+  bool build_BLAS_hair(Progress &progress,
+                       id<MTLDevice> device,
+                       id<MTLCommandQueue> queue,
+                       Geometry *const geom,
+                       bool refit);
+  bool build_BLAS_pointcloud(Progress &progress,
+                             id<MTLDevice> device,
+                             id<MTLCommandQueue> queue,
+                             Geometry *const geom,
+                             bool refit);
+  bool build_TLAS(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
+
+  API_AVAILABLE(macos(11.0))
+  void set_accel_struct(id<MTLAccelerationStructure> new_accel_struct);
+};
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/device.h

@@ -0,0 +1,30 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#include "util/string.h"
+#include "util/unique_ptr.h"
+#include "util/vector.h"
+
+CCL_NAMESPACE_BEGIN
+
+class Device;
+class DeviceInfo;
+class Profiler;
+class Stats;
+
+bool device_metal_init();
+void device_metal_exit();
+
+unique_ptr<Device> device_metal_create(const DeviceInfo &info,
+                                       Stats &stats,
+                                       Profiler &profiler,
+                                       bool headless);
+
+void device_metal_info(vector<DeviceInfo> &devices);
+
+string device_metal_capabilities();
+
+CCL_NAMESPACE_END

blender-5.2.0/intern/cycles/device/metal/device.mm

@@ -0,0 +1,155 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#include "device/device.h"
+
+#ifdef WITH_METAL
+
+#  include "device/metal/device.h"
+#  include "device/metal/device_impl.h"
+#  include "integrator/denoiser_oidn_gpu.h"
+
+#endif
+
+#include "util/debug.h"
+#include "util/set.h"
+#include "util/system.h"
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef WITH_METAL
+
+unique_ptr<Device> device_metal_create(const DeviceInfo &info,
+                                       Stats &stats,
+                                       Profiler &profiler,
+                                       bool headless)
+{
+  return make_unique<MetalDevice>(info, stats, profiler, headless);
+}
+
+bool device_metal_init()
+{
+  return true;
+}
+
+void device_metal_exit()
+{
+  MetalDeviceKernels::static_deinitialize();
+}
+
+void device_metal_info(vector<DeviceInfo> &devices)
+{
+  auto usable_devices = MetalInfo::get_usable_devices();
+  /* Devices are numbered consecutively across platforms. */
+  set<string> unique_ids;
+  int device_index = 0;
+  for (id<MTLDevice> &device : usable_devices) {
+    /* Compute unique ID for persistent user preferences. */
+    string device_name = MetalInfo::get_device_name(device);
+
+    string id = string("METAL_") + device_name;
+
+    /* Hardware ID might not be unique, add device number in that case. */
+    if (unique_ids.contains(id)) {
+      id += string_printf("_ID_%d", device_index);
+    }
+    unique_ids.insert(id);
+
+    /* Create DeviceInfo. */
+    DeviceInfo info;
+    info.type = DEVICE_METAL;
+    info.description = string_remove_trademark(string(device_name));
+
+    info.num = device_index;
+    /* We don't know if it's used for display, but assume it is. */
+    info.display_device = true;
+    info.denoisers = DENOISER_NONE;
+    info.id = id;
+#  if defined(WITH_OPENIMAGEDENOISE)
+#    if OIDN_VERSION >= 20300
+    if (oidnIsMetalDeviceSupported(device)) {
+#    else
+    if (OIDNDenoiserGPU::is_device_supported(info)) {
+#    endif
+      info.denoisers |= DENOISER_OPENIMAGEDENOISE;
+    }
+#  endif
+
+    info.has_nanovdb = true;
+
+    /* MNEE caused "Compute function exceeds available temporary registers" in macOS < 13 due to a
+     * bug in spill buffer allocation sizing. */
+    info.has_mnee_ = false;
+    if (@available(macos 13.0, *)) {
+      info.has_mnee_ = true;
+    }
+
+    info.use_hardware_raytracing = false;
+
+    /* MetalRT now uses features exposed in Xcode versions corresponding to macOS 14+, so don't
+     * expose it in builds from older Xcode versions. */
+#  if defined(MAC_OS_VERSION_14_0)
+    if (@available(macos 14.0, *)) {
+      info.use_hardware_raytracing = device.supportsRaytracing;
+
+      /* Use hardware raytracing for faster rendering on architectures that support it. */
+      info.use_metalrt_by_default = device.supportsRaytracing &&
+                                    (MetalInfo::get_apple_gpu_architecture(device) >= APPLE_M3);
+    }
+#  endif
+
+    devices.push_back(info);
+    device_index++;
+
+    LOG_INFO << "Added device \"" << info.description << "\" with id \"" << info.id << "\".";
+
+    if (info.denoisers & DENOISER_OPENIMAGEDENOISE) {
+      LOG_INFO << "Device with id \"" << info.id << "\" supports "
+               << denoiserTypeToHumanReadable(DENOISER_OPENIMAGEDENOISE) << ".";
+    }
+  }
+}
+
+string device_metal_capabilities()
+{
+  string result;
+  auto allDevices = MTLCopyAllDevices();
+  uint32_t num_devices = (uint32_t)allDevices.count;
+  if (num_devices == 0) {
+    return "No Metal devices found\n";
+  }
+  result += string_printf("Number of devices: %u\n", num_devices);
+
+  for (id<MTLDevice> device in allDevices) {
+    string device_name = MetalInfo::get_device_name(device);
+    result += string_printf("\t\tDevice: %s\n", device_name.c_str());
+  }
+
+  return result;
+}
+
+#else
+
+unique_ptr<Device> device_metal_create(const DeviceInfo & /*info*/,
+                                       Stats & /*stats*/,
+                                       Profiler & /*profiler*/)
+{
+  return nullptr;
+}
+
+bool device_metal_init()
+{
+  return false;
+}
+
+void device_metal_info(vector<DeviceInfo> & /*devices*/) {}
+
+string device_metal_capabilities()
+{
+  return "";
+}
+
+#endif
+
+CCL_NAMESPACE_END

blender-5.2.0/intern/cycles/device/metal/device_impl.h

@@ -0,0 +1,217 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#ifdef WITH_METAL
+
+#  include "bvh/bvh.h"
+#  include "device/device.h"
+#  include "device/metal/bvh.h"
+#  include "device/metal/device.h"
+#  include "device/metal/kernel.h"
+#  include "device/metal/queue.h"
+#  include "device/metal/util.h"
+
+#  include <Metal/Metal.h>
+
+CCL_NAMESPACE_BEGIN
+
+class DeviceQueue;
+
+class MetalDevice : public Device {
+ public:
+  id<MTLDevice> mtlDevice = nil;
+  id<MTLLibrary> mtlLibrary[PSO_NUM] = {nil};
+  id<MTLCommandQueue> mtlComputeCommandQueue = nil;
+  id<MTLCommandQueue> mtlGeneralCommandQueue = nil;
+  id<MTLCounterSampleBuffer> mtlCounterSampleBuffer = nil;
+  string source[PSO_NUM];
+  string kernels_md5[PSO_NUM];
+  string global_defines_md5[PSO_NUM];
+
+  bool capture_enabled = false;
+
+  /* Argument buffer for static data. */
+  id<MTLBuffer> launch_params_buffer = nil;
+  KernelParamsMetal *launch_params = nullptr;
+
+  /* MetalRT members ---------------------------------- */
+  bool use_metalrt = false;
+  bool use_metalrt_extended_limits = false;
+  bool motion_blur = false;
+  bool use_pcmi = false;
+
+  id<MTLBuffer> blas_buffer = nil;
+
+  API_AVAILABLE(macos(11.0))
+  vector<id<MTLAccelerationStructure>> unique_blas_array;
+
+  API_AVAILABLE(macos(11.0))
+  vector<id<MTLAccelerationStructure>> blas_array;
+
+  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);
+
+  /* For externally-owned resources (e.g. graphics interop buffers) which need to be resident for
+   * kernels to access them, but shouldn't be included in our stats. */
+  void add_to_residency_set(id<MTLResource> allocation);
+  void remove_from_residency_set(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;
+  int max_threads_per_threadgroup;
+
+  int mtlDevId = 0;
+  bool has_error = false;
+
+  struct MetalMem {
+    device_memory *mem = nullptr;
+    int pointer_index = -1;
+    id<MTLBuffer> mtlBuffer = nil;
+    id<MTLTexture> mtlTexture = nil;
+    uint64_t offset = 0;
+    uint64_t size = 0;
+    void *hostPtr = nullptr;
+  };
+  using MetalMemMap = map<device_memory *, unique_ptr<MetalMem>>;
+  MetalMemMap metal_mem_map;
+  std::vector<id<MTLResource>> delayed_free_list;
+  std::recursive_mutex metal_mem_map_mutex;
+
+  /* Bindless Textures */
+  bool is_texture(const KernelImageInfo &info);
+  device_vector<KernelImageInfo> image_info;
+  id<MTLBuffer> image_bindings = nil;
+  std::vector<id<MTLResource>> image_info_id_map;
+
+  MetalPipelineType kernel_specialization_level = PSO_GENERIC;
+
+  int device_id = 0;
+
+  static thread_mutex existing_devices_mutex;
+  static std::map<int, MetalDevice *> active_device_ids;
+
+  static bool is_device_cancelled(const int device_id);
+
+  static MetalDevice *get_device_by_ID(const int device_idID,
+                                       thread_scoped_lock &existing_devices_mutex_lock);
+
+  bool is_ready(string &status) const override;
+
+  void cancel() override;
+
+  BVHLayoutMask get_bvh_layout_mask(uint /*kernel_features*/) const override;
+
+  void set_error(const string &error) override;
+
+  MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler, bool headless);
+
+  ~MetalDevice() override;
+
+  bool support_device(const uint /*kernel_features*/);
+
+  bool check_peer_access(Device *peer_device) override;
+
+  bool use_adaptive_compilation();
+
+  bool use_local_atomic_sort() const;
+
+  string preprocess_source(MetalPipelineType pso_type,
+                           const uint kernel_features,
+                           string *source = nullptr);
+
+  void refresh_source_and_kernels_md5(MetalPipelineType pso_type);
+
+  void make_source(MetalPipelineType pso_type, const uint kernel_features);
+
+  bool load_kernels(const uint kernel_features) override;
+
+  void load_image_info();
+
+  void erase_allocation(device_memory &mem);
+
+  bool should_use_graphics_interop(const GraphicsInteropDevice &interop_device,
+                                   const bool log) override;
+
+  void *get_native_buffer(device_ptr ptr) override;
+
+  unique_ptr<DeviceQueue> gpu_queue_create() override;
+
+  void build_bvh(BVH *bvh, Progress &progress, bool refit) override;
+
+  bool set_bvh_limits(size_t instance_count, size_t max_prim_count) override;
+
+  void optimize_for_scene(Scene *scene) override;
+
+  static void compile_and_load(const int device_id, MetalPipelineType pso_type);
+
+  /* ------------------------------------------------------------------ */
+  /* low-level memory management */
+
+  bool max_working_set_exceeded(const size_t safety_margin = 8 * 1024 * 1024) const;
+
+  MetalMem *generic_alloc(device_memory &mem);
+
+  void generic_copy_to(device_memory &mem);
+
+  void generic_free(device_memory &mem);
+
+  void mem_alloc(device_memory &mem) override;
+
+  void mem_copy_to(device_memory &mem) override;
+
+  void mem_move_to_host(device_memory &mem) override;
+
+  void mem_copy_from(device_memory &mem)
+  {
+    mem_copy_from(mem, -1, -1, -1, -1);
+  }
+  void mem_copy_from(
+      device_memory &mem, const size_t y, size_t w, const size_t h, size_t elem) override;
+
+  void mem_or_from_device(device_memory &mem) override;
+
+  void mem_zero(device_memory &mem) override;
+
+  void mem_free(device_memory &mem) override;
+
+  device_ptr mem_alloc_sub_ptr(device_memory &mem, const size_t offset, size_t /*size*/) override;
+
+  void const_copy_to(const char *name, void *host, const size_t size) override;
+
+  void global_alloc(device_memory &mem);
+  void global_free(device_memory &mem);
+
+  void image_alloc(device_image &mem);
+  void image_alloc_as_buffer(device_image &mem);
+  void image_copy_to(device_image &mem);
+  void image_free(device_image &mem);
+
+  bool has_unified_memory() const override;
+
+  void flush_delayed_free_list();
+
+  void free_bvh();
+
+  void update_bvh(BVHMetal *bvh_metal);
+};
+
+CCL_NAMESPACE_END
+
+#endif

blender-5.2.0/intern/cycles/device/metal/graphics_interop.h

@@ -0,0 +1,48 @@
+/* SPDX-FileCopyrightText: 2025 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#ifdef WITH_METAL
+
+#  include "device/graphics_interop.h"
+#  include "device/metal/device_impl.h"
+
+#  include "session/display_driver.h"
+
+CCL_NAMESPACE_BEGIN
+
+class MetalDevice;
+class MetalDeviceQueue;
+
+class MetalDeviceGraphicsInterop : public DeviceGraphicsInterop {
+ public:
+  explicit MetalDeviceGraphicsInterop(MetalDeviceQueue *queue);
+
+  MetalDeviceGraphicsInterop(const MetalDeviceGraphicsInterop &other) = delete;
+  MetalDeviceGraphicsInterop(MetalDeviceGraphicsInterop &&other) noexcept = delete;
+
+  ~MetalDeviceGraphicsInterop() override;
+
+  MetalDeviceGraphicsInterop &operator=(const MetalDeviceGraphicsInterop &other) = delete;
+  MetalDeviceGraphicsInterop &operator=(MetalDeviceGraphicsInterop &&other) = delete;
+
+  void set_buffer(GraphicsInteropBuffer &interop_buffer) override;
+
+  device_ptr map() override;
+  void unmap() override;
+
+ protected:
+  MetalDeviceQueue *queue_ = nullptr;
+  MetalDevice *device_ = nullptr;
+
+  /* Native handle. */
+  MetalDevice::MetalMem mem_;
+  size_t size_ = 0;
+
+  /* The destination was requested to be cleared. */
+  bool need_zero_ = false;
+};
+
+CCL_NAMESPACE_END
+
+#endif

blender-5.2.0/intern/cycles/device/metal/graphics_interop.mm

@@ -0,0 +1,60 @@
+/* SPDX-FileCopyrightText: 2025 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#ifdef WITH_METAL
+
+#  include "device/metal/graphics_interop.h"
+
+#  include "device/metal/device_impl.h"
+
+CCL_NAMESPACE_BEGIN
+
+MetalDeviceGraphicsInterop::MetalDeviceGraphicsInterop(MetalDeviceQueue *queue)
+    : queue_(queue), device_(static_cast<MetalDevice *>(queue->device))
+{
+}
+
+MetalDeviceGraphicsInterop::~MetalDeviceGraphicsInterop()
+{
+  device_->remove_from_residency_set(mem_.mtlBuffer);
+}
+
+void MetalDeviceGraphicsInterop::set_buffer(GraphicsInteropBuffer &interop_buffer)
+{
+  if (interop_buffer.is_empty()) {
+    device_->remove_from_residency_set(mem_.mtlBuffer);
+    mem_.mtlBuffer = nullptr;
+    size_ = 0;
+    return;
+  }
+
+  need_zero_ |= interop_buffer.take_zero();
+
+  if (!interop_buffer.has_new_handle()) {
+    return;
+  }
+
+  /* The interop buffer is externally owned, so it doesn't go through metal_mem_alloc. We still
+   * need to put it in the residency set so kernels can write display output into it. */
+  device_->remove_from_residency_set(mem_.mtlBuffer);
+  mem_.mtlBuffer = reinterpret_cast<id<MTLBuffer>>(interop_buffer.take_handle());
+  size_ = interop_buffer.get_size();
+  device_->add_to_residency_set(mem_.mtlBuffer);
+}
+
+device_ptr MetalDeviceGraphicsInterop::map()
+{
+  if (mem_.mtlBuffer && need_zero_) {
+    memset([mem_.mtlBuffer contents], 0, size_);
+    need_zero_ = false;
+  }
+
+  return device_ptr(&mem_);
+}
+
+void MetalDeviceGraphicsInterop::unmap() {}
+
+CCL_NAMESPACE_END
+
+#endif

blender-5.2.0/intern/cycles/device/metal/kernel.h

@@ -0,0 +1,134 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#ifdef WITH_METAL
+
+#  include "device/kernel.h"
+
+#  include <Metal/Metal.h>
+
+CCL_NAMESPACE_BEGIN
+
+class MetalDevice;
+
+enum {
+  METALRT_TABLE_DEFAULT,
+  METALRT_TABLE_SHADOW,
+  METALRT_TABLE_SHADOW_ALL,
+  METALRT_TABLE_VOLUME,
+  METALRT_TABLE_LOCAL,
+  METALRT_TABLE_LOCAL_MBLUR,
+  METALRT_TABLE_LOCAL_SINGLE_HIT,
+  METALRT_TABLE_LOCAL_SINGLE_HIT_MBLUR,
+  METALRT_TABLE_NUM
+};
+
+/* Pipeline State Object types */
+enum MetalPipelineType {
+  /* A kernel that can be used with all scenes, supporting all features.
+   * It is slow to compile, but only needs to be compiled once and is then
+   * cached for future render sessions. This allows a render to get underway
+   * on the GPU quickly.
+   */
+  PSO_GENERIC,
+
+  /* A intersection kernel that is very quick to specialize and results in faster intersection
+   * kernel performance. It uses Metal function constants to replace several KernelData variables
+   * with fixed constants.
+   */
+  PSO_SPECIALIZED_INTERSECT,
+
+  /* A shading kernel that is slow to specialize, but results in faster shading kernel performance
+   * rendered. It uses Metal function constants to replace several KernelData variables with fixed
+   * constants and short-circuit all unused SVM node case handlers.
+   */
+  PSO_SPECIALIZED_SHADE,
+
+  PSO_NUM
+};
+
+#  define METALRT_FEATURE_MASK \
+    (KERNEL_FEATURE_HAIR | KERNEL_FEATURE_HAIR_THICK | KERNEL_FEATURE_POINTCLOUD)
+
+const char *kernel_type_as_string(MetalPipelineType pso_type);
+
+/* A pipeline object that can be shared between multiple instances of MetalDeviceQueue. */
+class MetalKernelPipeline {
+ public:
+  void compile();
+
+  int pipeline_id;
+  int originating_device_id;
+
+  id<MTLLibrary> mtlLibrary = nil;
+  MetalPipelineType pso_type;
+  string kernels_md5;
+  size_t usage_count = 0;
+
+  KernelData kernel_data_;
+  bool use_metalrt;
+  uint32_t kernel_features = 0;
+
+  int threads_per_threadgroup;
+
+  DeviceKernel device_kernel;
+  bool loaded = false;
+  id<MTLDevice> mtlDevice = nil;
+  id<MTLFunction> function = nil;
+  id<MTLComputePipelineState> pipeline = nil;
+  int num_threads_per_block = 0;
+
+  bool should_use_binary_archive() const;
+  id<MTLFunction> make_intersection_function(const char *function_name);
+
+  string error_str;
+
+  NSArray *table_functions[METALRT_TABLE_NUM] = {nil};
+};
+
+/* An actively instanced pipeline that can only be used by a single instance of MetalDeviceQueue.
+ */
+class MetalDispatchPipeline {
+ public:
+  ~MetalDispatchPipeline();
+
+  bool update(MetalDevice *metal_device, DeviceKernel kernel);
+  void free_intersection_function_tables();
+
+ private:
+  friend class MetalDeviceQueue;
+  friend struct ShaderCache;
+
+  int pipeline_id = -1;
+
+  MetalDevice *metal_device = nullptr;
+  MetalPipelineType pso_type;
+  id<MTLComputePipelineState> pipeline = nil;
+  int num_threads_per_block = 0;
+
+  API_AVAILABLE(macos(11.0))
+  id<MTLIntersectionFunctionTable> intersection_func_table[METALRT_TABLE_NUM] = {nil};
+};
+
+/* Cache of Metal kernels for each DeviceKernel. */
+namespace MetalDeviceKernels {
+
+int num_incomplete_specialization_requests();
+int get_loaded_kernel_count(const MetalDevice *device, MetalPipelineType pso_type);
+bool should_load_kernels(const MetalDevice *device, MetalPipelineType pso_type);
+bool load(MetalDevice *device, MetalPipelineType pso_type);
+const MetalKernelPipeline *get_best_pipeline(const MetalDevice *device, DeviceKernel kernel);
+void wait_for_all();
+bool is_benchmark_warmup();
+
+/* Deinitialize all static variables, so that no code would run on application exit. */
+void static_deinitialize();
+
+} /* namespace MetalDeviceKernels */
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/kernel.mm

@@ -0,0 +1,935 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#ifdef WITH_METAL
+
+#  include <algorithm>
+#  include <atomic>
+#  include <chrono>
+#  include <deque>
+#  include <thread>
+#  include <vector>
+
+#  include "device/metal/device_impl.h"
+#  include "device/metal/kernel.h"
+
+#  include "kernel/device/metal/function_constants.h"
+
+#  include "util/debug.h"
+#  include "util/md5.h"
+#  include "util/path.h"
+#  include "util/tbb.h"
+#  include "util/time.h"
+#  include "util/unique_ptr.h"
+
+CCL_NAMESPACE_BEGIN
+
+const char *kernel_type_as_string(MetalPipelineType pso_type)
+{
+  switch (pso_type) {
+    case PSO_GENERIC:
+      return "PSO_GENERIC";
+    case PSO_SPECIALIZED_INTERSECT:
+      return "PSO_SPECIALIZED_INTERSECT";
+    case PSO_SPECIALIZED_SHADE:
+      return "PSO_SPECIALIZED_SHADE";
+    default:
+      assert(0);
+  }
+  return "";
+}
+
+struct ShaderCache {
+  ShaderCache(id<MTLDevice> _mtlDevice) : mtlDevice(_mtlDevice)
+  {
+    /* Initialize occupancy tuning LUT. */
+
+    // TODO: Look into tuning for DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT and
+    // DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_*.
+
+    switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
+      default:
+      case APPLE_M3:
+        /* Peak occupancy is achieved through Dynamic Caching on M3 GPUs. */
+        for (size_t i = 0; i < DEVICE_KERNEL_NUM; i++) {
+          occupancy_tuning[i] = {64, 64};
+        }
+        break;
+      case APPLE_M2_BIG:
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {384, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {640, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST] = {1024, 64};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] = {704, 704};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE] = {640, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY] = {896, 768};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND] = {512, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] = {32, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = {768, 576};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY] = {896, 768};
+        break;
+      case APPLE_M2:
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {32, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {832, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST] = {64, 64};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] = {64, 64};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE] = {704, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY] = {1024, 256};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND] = {64, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] = {256, 256};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = {448, 384};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY] = {1024, 1024};
+        break;
+      case APPLE_M1:
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {256, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {768, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST] = {512, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] = {384, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE] = {512, 64};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY] = {512, 256};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND] = {512, 128};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] = {384, 32};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = {576, 384};
+        occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY] = {832, 832};
+        break;
+    }
+
+    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS] = {1024, 1024};
+    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS] = {1024, 1024};
+  }
+  ~ShaderCache();
+
+  /* Get the fastest available pipeline for the specified kernel. */
+  MetalKernelPipeline *get_best_pipeline(DeviceKernel kernel, const MetalDevice *device);
+
+  /* Non-blocking request for a kernel, optionally specialized to the scene being rendered by
+   * device. */
+  void load_kernel(DeviceKernel kernel, MetalDevice *device, MetalPipelineType pso_type);
+
+  bool should_load_kernel(DeviceKernel device_kernel,
+                          const MetalDevice *device,
+                          MetalPipelineType pso_type);
+
+  void wait_for_all();
+
+  friend ShaderCache *get_shader_cache(id<MTLDevice> mtlDevice);
+
+  void compile_thread_func();
+
+  using PipelineCollection = std::vector<unique_ptr<MetalKernelPipeline>>;
+
+  struct OccupancyTuningParameters {
+    int threads_per_threadgroup = 0;
+    int num_threads_per_block = 0;
+  } occupancy_tuning[DEVICE_KERNEL_NUM];
+
+  std::mutex cache_mutex;
+
+  PipelineCollection pipelines[DEVICE_KERNEL_NUM];
+  id<MTLDevice> mtlDevice;
+
+  static bool running;
+  std::condition_variable cond_var;
+  std::deque<unique_ptr<MetalKernelPipeline>> request_queue;
+  std::vector<std::thread> compile_threads;
+  std::atomic_int incomplete_requests = 0;
+  std::atomic_int incomplete_specialization_requests = 0;
+};
+
+bool ShaderCache::running = true;
+
+const int MAX_POSSIBLE_GPUS_ON_SYSTEM = 8;
+using DeviceShaderCache = std::pair<id<MTLDevice>, unique_ptr<ShaderCache>>;
+int g_shaderCacheCount = 0;
+DeviceShaderCache g_shaderCache[MAX_POSSIBLE_GPUS_ON_SYSTEM];
+
+/* Next UID for associating a MetalDispatchPipeline with an originating MetalKernelPipeline. */
+static std::atomic_int g_next_pipeline_id = 0;
+
+ShaderCache *get_shader_cache(id<MTLDevice> mtlDevice)
+{
+  for (int i = 0; i < g_shaderCacheCount; i++) {
+    if (g_shaderCache[i].first == mtlDevice) {
+      return g_shaderCache[i].second.get();
+    }
+  }
+
+  static thread_mutex g_shaderCacheCountMutex;
+  g_shaderCacheCountMutex.lock();
+  int index = g_shaderCacheCount++;
+  g_shaderCacheCountMutex.unlock();
+
+  assert(index < MAX_POSSIBLE_GPUS_ON_SYSTEM);
+  g_shaderCache[index].first = mtlDevice;
+  g_shaderCache[index].second = make_unique<ShaderCache>(mtlDevice);
+  return g_shaderCache[index].second.get();
+}
+
+ShaderCache::~ShaderCache()
+{
+  running = false;
+  cond_var.notify_all();
+
+  metal_printf("Waiting for ShaderCache threads... (incomplete_requests = %d)",
+               int(incomplete_requests));
+  for (auto &thread : compile_threads) {
+    thread.join();
+  }
+  metal_printf("ShaderCache shut down.");
+}
+
+void ShaderCache::wait_for_all()
+{
+  while (incomplete_requests > 0) {
+    std::this_thread::sleep_for(std::chrono::milliseconds(100));
+  }
+}
+
+void ShaderCache::compile_thread_func()
+{
+  while (running) {
+
+    /* wait for / acquire next request */
+    unique_ptr<MetalKernelPipeline> pipeline;
+    {
+      thread_scoped_lock lock(cache_mutex);
+      cond_var.wait(lock, [&] { return !running || !request_queue.empty(); });
+      if (!running || request_queue.empty()) {
+        continue;
+      }
+
+      pipeline = std::move(request_queue.front());
+      request_queue.pop_front();
+    }
+
+    /* Service the request. */
+    DeviceKernel device_kernel = pipeline->device_kernel;
+    MetalPipelineType pso_type = pipeline->pso_type;
+
+    if (MetalDevice::is_device_cancelled(pipeline->originating_device_id)) {
+      /* The originating MetalDevice is no longer active, so this request is obsolete. */
+      metal_printf("Cancelling compilation of %s (%s)",
+                   device_kernel_as_string(device_kernel),
+                   kernel_type_as_string(pso_type));
+    }
+    else {
+      /* Do the actual compilation. */
+      pipeline->compile();
+
+      thread_scoped_lock lock(cache_mutex);
+      auto &collection = pipelines[device_kernel];
+
+      /* Cache up to 3 kernel variants with the same pso_type in memory, purging oldest first. */
+      int max_entries_of_same_pso_type = 3;
+      for (int i = (int)collection.size() - 1; i >= 0; i--) {
+        if (collection[i]->pso_type == pso_type) {
+          max_entries_of_same_pso_type -= 1;
+          if (max_entries_of_same_pso_type == 0) {
+            metal_printf("Purging oldest %s:%s kernel from ShaderCache",
+                         kernel_type_as_string(pso_type),
+                         device_kernel_as_string(device_kernel));
+            collection.erase(collection.begin() + i);
+            break;
+          }
+        }
+      }
+      collection.push_back(std::move(pipeline));
+    }
+    incomplete_requests--;
+    if (pso_type != PSO_GENERIC) {
+      incomplete_specialization_requests--;
+    }
+  }
+}
+
+bool ShaderCache::should_load_kernel(DeviceKernel device_kernel,
+                                     const MetalDevice *device,
+                                     MetalPipelineType pso_type)
+{
+  if (!running) {
+    return false;
+  }
+
+  if (!device_kernel_has_gpu_function(device_kernel)) {
+    /* Skip megakernel and other markers without a GPU function. */
+    return false;
+  }
+
+  if (device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE) {
+    if ((device->kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) == 0) {
+      /* Skip shade_surface_raytrace kernel if the scene doesn't require it. */
+      return false;
+    }
+  }
+
+  if (device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE) {
+    if ((device->kernel_features & KERNEL_FEATURE_MNEE) == 0) {
+      /* Skip the MNEE kernel if the scene doesn't require it. */
+      return false;
+    }
+  }
+
+  if (pso_type != PSO_GENERIC) {
+    /* Only specialize kernels where it can make an impact. */
+    if (device_kernel < DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST ||
+        device_kernel > DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL)
+    {
+      return false;
+    }
+
+    /* Only specialize shading / intersection kernels as requested. */
+    bool is_shade_kernel = (device_kernel >= DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
+    bool is_shade_pso = (pso_type == PSO_SPECIALIZED_SHADE);
+    if (is_shade_pso != is_shade_kernel) {
+      return false;
+    }
+  }
+
+  {
+    /* check whether the kernel has already been requested / cached */
+    thread_scoped_lock lock(cache_mutex);
+    for (auto &pipeline : pipelines[device_kernel]) {
+      if (pipeline->kernels_md5 == device->kernels_md5[pso_type]) {
+        return false;
+      }
+    }
+  }
+
+  return true;
+}
+
+void ShaderCache::load_kernel(DeviceKernel device_kernel,
+                              MetalDevice *device,
+                              MetalPipelineType pso_type)
+{
+  {
+    /* create compiler threads on first run */
+    thread_scoped_lock lock(cache_mutex);
+    if (compile_threads.empty()) {
+      /* Limit to 2 MTLCompiler instances by default. In macOS >= 13.3 we can query the upper
+       * limit. */
+      int max_mtlcompiler_threads = 2;
+
+#  if defined(MAC_OS_VERSION_13_3)
+      if (@available(macOS 13.3, *)) {
+        /* Subtract one to avoid contention with the real-time GPU module. */
+        max_mtlcompiler_threads = max(2,
+                                      int([mtlDevice maximumConcurrentCompilationTaskCount]) - 1);
+      }
+#  endif
+
+      metal_printf("Spawning %d Cycles kernel compilation threads", max_mtlcompiler_threads);
+      for (int i = 0; i < max_mtlcompiler_threads; i++) {
+        compile_threads.emplace_back([this] { this->compile_thread_func(); });
+      }
+    }
+  }
+
+  if (!should_load_kernel(device_kernel, device, pso_type)) {
+    return;
+  }
+
+  incomplete_requests++;
+  if (pso_type != PSO_GENERIC) {
+    incomplete_specialization_requests++;
+  }
+
+  unique_ptr<MetalKernelPipeline> pipeline = make_unique<MetalKernelPipeline>();
+
+  /* Keep track of the originating device's ID so that we can cancel requests if the device ceases
+   * to be active. */
+  pipeline->pipeline_id = g_next_pipeline_id.fetch_add(1);
+  pipeline->originating_device_id = device->device_id;
+  pipeline->kernel_data_ = device->launch_params->data;
+  pipeline->pso_type = pso_type;
+  pipeline->mtlDevice = mtlDevice;
+  pipeline->kernels_md5 = device->kernels_md5[pso_type];
+  pipeline->mtlLibrary = device->mtlLibrary[pso_type];
+  pipeline->device_kernel = device_kernel;
+  pipeline->threads_per_threadgroup = device->max_threads_per_threadgroup;
+
+  if (occupancy_tuning[device_kernel].threads_per_threadgroup) {
+    pipeline->threads_per_threadgroup = occupancy_tuning[device_kernel].threads_per_threadgroup;
+    pipeline->num_threads_per_block = occupancy_tuning[device_kernel].num_threads_per_block;
+  }
+
+  /* metalrt options */
+  pipeline->use_metalrt = device->use_metalrt;
+  pipeline->kernel_features = device->kernel_features;
+
+  {
+    thread_scoped_lock lock(cache_mutex);
+    request_queue.push_back(std::move(pipeline));
+  }
+  cond_var.notify_one();
+}
+
+MetalKernelPipeline *ShaderCache::get_best_pipeline(DeviceKernel kernel, const MetalDevice *device)
+{
+  while (running && !device->has_error) {
+    /* Search all loaded pipelines with matching kernels_md5 checksums. */
+    MetalKernelPipeline *best_match = nullptr;
+    {
+      thread_scoped_lock lock(cache_mutex);
+      for (auto &candidate : pipelines[kernel]) {
+        if (candidate->loaded &&
+            candidate->kernels_md5 == device->kernels_md5[candidate->pso_type])
+        {
+          /* Replace existing match if candidate is more specialized. */
+          if (!best_match || candidate->pso_type > best_match->pso_type) {
+            best_match = candidate.get();
+          }
+        }
+      }
+    }
+
+    if (best_match) {
+      if (best_match->usage_count == 0 && best_match->pso_type != PSO_GENERIC) {
+        metal_printf("Swapping in %s version of %s",
+                     kernel_type_as_string(best_match->pso_type),
+                     device_kernel_as_string(kernel));
+      }
+      best_match->usage_count += 1;
+      return best_match;
+    }
+
+    /* Spin until a matching kernel is loaded, or we're shutting down. */
+    std::this_thread::sleep_for(std::chrono::milliseconds(100));
+  }
+  return nullptr;
+}
+
+bool MetalKernelPipeline::should_use_binary_archive() const
+{
+  /* Issues with binary archives in older macOS versions. */
+  if (@available(macOS 15.4, *)) {
+    if (auto *str = getenv("CYCLES_METAL_DISABLE_BINARY_ARCHIVES")) {
+      if (atoi(str) != 0) {
+        /* Don't archive if we have opted out by env var. */
+        return false;
+      }
+    }
+
+    if (use_metalrt && device_kernel_has_intersection(device_kernel)) {
+      /* Binary linked functions aren't supported in binary archives. */
+      return false;
+    }
+
+    if (pso_type == PSO_GENERIC) {
+      /* Archive the generic kernels. */
+      return true;
+    }
+
+    if ((device_kernel >= DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND &&
+         device_kernel <= DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW) ||
+        (device_kernel >= DEVICE_KERNEL_SHADER_EVAL_DISPLACE &&
+         device_kernel <= DEVICE_KERNEL_SHADER_EVAL_VOLUME_DENSITY))
+    {
+      /* Archive all shade kernels - they take a long time to compile. */
+      return true;
+    }
+
+    /* The remaining kernels are all fast to compile. They may get cached by the system shader
+     * cache, but will be quick to regenerate if not. */
+  }
+  return false;
+}
+
+static MTLFunctionConstantValues *GetConstantValues(const KernelData *data = nullptr)
+{
+  MTLFunctionConstantValues *constant_values = [MTLFunctionConstantValues new];
+
+  MTLDataType MTLDataType_int = MTLDataTypeInt;
+  MTLDataType MTLDataType_float = MTLDataTypeFloat;
+  MTLDataType MTLDataType_float2 = MTLDataTypeFloat2;
+  MTLDataType MTLDataType_float4 = MTLDataTypeFloat4;
+  KernelData zero_data = {0};
+  if (!data) {
+    data = &zero_data;
+  }
+  [constant_values setConstantValue:&zero_data type:MTLDataType_int atIndex:Kernel_DummyConstant];
+
+  bool next_member_is_specialized = true;
+
+#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
+
+#  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
+    [constant_values setConstantValue:next_member_is_specialized ? (void *)&data->parent.name : \
+                                                                   (void *)&zero_data \
+                                 type:MTLDataType_##_type \
+                              atIndex:KernelData_##parent##_##name]; \
+    next_member_is_specialized = true;
+
+#  include "kernel/data_template.h"
+
+  [constant_values setConstantValue:&data->kernel_features
+                               type:MTLDataTypeInt
+                            atIndex:KernelData_kernel_features];
+
+  return constant_values;
+}
+
+void MetalDispatchPipeline::free_intersection_function_tables()
+{
+  for (int table = 0; table < METALRT_TABLE_NUM; table++) {
+    if (intersection_func_table[table]) {
+      /* 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;
+    }
+  }
+}
+
+MetalDispatchPipeline::~MetalDispatchPipeline()
+{
+  free_intersection_function_tables();
+}
+
+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) {
+    return false;
+  }
+
+  if (pipeline_id == best_pipeline->pipeline_id) {
+    /* The best pipeline is already active - nothing to do. */
+    return true;
+  }
+  pipeline_id = best_pipeline->pipeline_id;
+  pipeline = best_pipeline->pipeline;
+  pso_type = best_pipeline->pso_type;
+  num_threads_per_block = best_pipeline->num_threads_per_block;
+
+  /* Create the MTLIntersectionFunctionTables if needed. */
+  if (best_pipeline->use_metalrt && device_kernel_has_intersection(best_pipeline->device_kernel)) {
+    free_intersection_function_tables();
+
+    for (int table = 0; table < METALRT_TABLE_NUM; table++) {
+      @autoreleasepool {
+        MTLIntersectionFunctionTableDescriptor *ift_desc =
+            [[MTLIntersectionFunctionTableDescriptor alloc] init];
+        ift_desc.functionCount = best_pipeline->table_functions[table].count;
+        intersection_func_table[table] = [this->pipeline
+            newIntersectionFunctionTableWithDescriptor:ift_desc];
+
+        /* Finally write the function handles into this pipeline's table */
+        int size = int([best_pipeline->table_functions[table] count]);
+        for (int i = 0; i < size; i++) {
+          id<MTLFunctionHandle> handle = [pipeline
+              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]);
+      }
+    }
+  }
+
+  return true;
+}
+
+id<MTLFunction> MetalKernelPipeline::make_intersection_function(const char *function_name)
+{
+  MTLFunctionDescriptor *desc = [MTLIntersectionFunctionDescriptor functionDescriptor];
+  desc.name = [@(function_name) copy];
+
+  if (pso_type != PSO_GENERIC) {
+    desc.constantValues = GetConstantValues(&kernel_data_);
+  }
+  else {
+    desc.constantValues = GetConstantValues();
+  }
+
+  NSError *error = nullptr;
+  id<MTLFunction> rt_intersection_function = [mtlLibrary newFunctionWithDescriptor:desc
+                                                                             error:&error];
+
+  if (rt_intersection_function == nil) {
+    NSString *err = [error localizedDescription];
+    string errors = [err UTF8String];
+
+    error_str = string_printf(
+        "Error getting intersection function \"%s\": %s", function_name, errors.c_str());
+  }
+  else {
+    rt_intersection_function.label = [@(function_name) copy];
+  }
+  return rt_intersection_function;
+}
+
+void MetalKernelPipeline::compile()
+{
+  const std::string function_name = std::string("cycles_metal_") +
+                                    device_kernel_as_string(device_kernel);
+
+  NSError *error = nullptr;
+
+  MTLFunctionDescriptor *func_desc = [MTLIntersectionFunctionDescriptor functionDescriptor];
+  func_desc.name = [@(function_name.c_str()) copy];
+
+  if (pso_type != PSO_GENERIC) {
+    func_desc.constantValues = GetConstantValues(&kernel_data_);
+  }
+  else {
+    func_desc.constantValues = GetConstantValues();
+  }
+
+  function = [mtlLibrary newFunctionWithDescriptor:func_desc error:&error];
+
+  if (function == nil) {
+    NSString *err = [error localizedDescription];
+    string errors = [err UTF8String];
+    metal_printf("Error getting function \"%s\": %s", function_name.c_str(), errors.c_str());
+    return;
+  }
+
+  function.label = [@(function_name.c_str()) copy];
+
+  NSArray *linked_functions = nil;
+
+  if (use_metalrt && device_kernel_has_intersection(device_kernel)) {
+
+    NSMutableSet *unique_functions = [[NSMutableSet alloc] init];
+
+    auto add_intersection_functions = [&](int table_index,
+                                          const char *tri_fn,
+                                          const char *curve_fn = nullptr,
+                                          const char *point_fn = nullptr) {
+      table_functions[table_index] = [NSArray
+          arrayWithObjects:make_intersection_function(tri_fn),
+                           curve_fn ? make_intersection_function(curve_fn) : nil,
+                           point_fn ? make_intersection_function(point_fn) : nil,
+                           nil];
+
+      [unique_functions addObjectsFromArray:table_functions[table_index]];
+    };
+
+    add_intersection_functions(METALRT_TABLE_DEFAULT,
+                               "__intersection__tri",
+                               "__intersection__curve",
+                               "__intersection__point");
+    add_intersection_functions(METALRT_TABLE_SHADOW,
+                               "__intersection__tri_shadow",
+                               "__intersection__curve_shadow",
+                               "__intersection__point_shadow");
+    add_intersection_functions(METALRT_TABLE_SHADOW_ALL,
+                               "__intersection__tri_shadow_all",
+                               "__intersection__curve_shadow_all",
+                               "__intersection__point_shadow_all");
+    add_intersection_functions(METALRT_TABLE_VOLUME, "__intersection__volume_tri");
+    add_intersection_functions(METALRT_TABLE_LOCAL, "__intersection__local_tri");
+    add_intersection_functions(METALRT_TABLE_LOCAL_MBLUR, "__intersection__local_tri_mblur");
+    add_intersection_functions(METALRT_TABLE_LOCAL_SINGLE_HIT,
+                               "__intersection__local_tri_single_hit");
+    add_intersection_functions(METALRT_TABLE_LOCAL_SINGLE_HIT_MBLUR,
+                               "__intersection__local_tri_single_hit_mblur");
+
+    linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]]
+        sortedArrayUsingComparator:^NSComparisonResult(id<MTLFunction> f1, id<MTLFunction> f2) {
+          return [f1.label compare:f2.label];
+        }];
+    unique_functions = nil;
+  }
+
+  MTLComputePipelineDescriptor *computePipelineStateDescriptor =
+      [[MTLComputePipelineDescriptor alloc] init];
+
+  computePipelineStateDescriptor.buffers[0].mutability = MTLMutabilityImmutable;
+  computePipelineStateDescriptor.buffers[1].mutability = MTLMutabilityImmutable;
+  computePipelineStateDescriptor.buffers[2].mutability = MTLMutabilityImmutable;
+
+  computePipelineStateDescriptor.maxTotalThreadsPerThreadgroup = threads_per_threadgroup;
+  computePipelineStateDescriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth = true;
+
+  computePipelineStateDescriptor.computeFunction = function;
+
+  /* Attach the additional functions to an MTLLinkedFunctions object */
+  if (linked_functions) {
+    computePipelineStateDescriptor.linkedFunctions = [[MTLLinkedFunctions alloc] init];
+    computePipelineStateDescriptor.linkedFunctions.functions = linked_functions;
+  }
+  computePipelineStateDescriptor.maxCallStackDepth = 1;
+  if (use_metalrt && device_kernel_has_intersection(device_kernel)) {
+    computePipelineStateDescriptor.maxCallStackDepth = 2;
+  }
+
+  MTLPipelineOption pipelineOptions = MTLPipelineOptionNone;
+
+  bool use_binary_archive = should_use_binary_archive();
+  bool loading_existing_archive = false;
+  bool creating_new_archive = false;
+
+  id<MTLBinaryArchive> archive = nil;
+  string metalbin_path;
+  string metalbin_name;
+  if (use_binary_archive) {
+    NSProcessInfo *processInfo = [NSProcessInfo processInfo];
+    string osVersion = [[processInfo operatingSystemVersionString] UTF8String];
+    MD5Hash local_md5;
+    local_md5.append(kernels_md5);
+    local_md5.append(osVersion);
+    local_md5.append((uint8_t *)&this->threads_per_threadgroup,
+                     sizeof(this->threads_per_threadgroup));
+
+    /* Replace non-alphanumerical characters with underscores. */
+    string device_name = [mtlDevice.name UTF8String];
+    for (char &c : device_name) {
+      if ((c < '0' || c > '9') && (c < 'a' || c > 'z') && (c < 'A' || c > 'Z')) {
+        c = '_';
+      }
+    }
+
+    metalbin_name = device_name;
+    metalbin_name = path_join(metalbin_name, device_kernel_as_string(device_kernel));
+    metalbin_name = path_join(metalbin_name, kernel_type_as_string(pso_type));
+    metalbin_name = path_join(metalbin_name, local_md5.get_hex() + ".bin");
+
+    metalbin_path = path_cache_get(path_join("kernels", metalbin_name));
+    path_create_directories(metalbin_path);
+
+    /* Check if shader binary exists on disk, and if so, update the file timestamp for LRU purging
+     * to work as intended. */
+    loading_existing_archive = path_cache_kernel_exists_and_mark_used(metalbin_path);
+    creating_new_archive = !loading_existing_archive;
+
+    MTLBinaryArchiveDescriptor *archiveDesc = [[MTLBinaryArchiveDescriptor alloc] init];
+    if (loading_existing_archive) {
+      archiveDesc.url = [NSURL fileURLWithPath:@(metalbin_path.c_str())];
+    }
+    NSError *error = nil;
+    archive = [mtlDevice newBinaryArchiveWithDescriptor:archiveDesc error:&error];
+    if (!archive) {
+      const char *err = error ? [[error localizedDescription] UTF8String] : nullptr;
+      metal_printf("newBinaryArchiveWithDescriptor failed: %s", err ? err : "nil");
+    }
+    [archiveDesc release];
+
+    if (loading_existing_archive) {
+      pipelineOptions = MTLPipelineOptionFailOnBinaryArchiveMiss;
+      computePipelineStateDescriptor.binaryArchives = [NSArray arrayWithObjects:archive, nil];
+    }
+  }
+
+  bool recreate_archive = false;
+
+  /* Lambda to do the actual pipeline compilation. */
+  auto do_compilation = [&]() {
+    __block bool compilation_finished = false;
+    __block string error_str;
+
+    if (loading_existing_archive || !DebugFlags().metal.use_async_pso_creation) {
+      /* Use the blocking variant of newComputePipelineStateWithDescriptor if an archive exists on
+       * disk. It should load almost instantaneously, and will fail gracefully when loading a
+       * corrupt archive (unlike the async variant). */
+      NSError *error = nil;
+      pipeline = [mtlDevice newComputePipelineStateWithDescriptor:computePipelineStateDescriptor
+                                                          options:pipelineOptions
+                                                       reflection:nullptr
+                                                            error:&error];
+      const char *err = error ? [[error localizedDescription] UTF8String] : nullptr;
+      error_str = err ? err : "nil";
+    }
+    else {
+      /* Use the async variant of newComputePipelineStateWithDescriptor if no archive exists on
+       * disk. This allows us to respond to app shutdown. */
+      [mtlDevice
+          newComputePipelineStateWithDescriptor:computePipelineStateDescriptor
+                                        options:pipelineOptions
+                              completionHandler:^(id<MTLComputePipelineState> computePipelineState,
+                                                  MTLComputePipelineReflection * /*reflection*/,
+                                                  NSError *error) {
+                                pipeline = computePipelineState;
+
+                                /* Retain the pipeline so we can use it safely past the completion
+                                 * handler. */
+                                if (pipeline) {
+                                  [pipeline retain];
+                                }
+                                const char *err = error ?
+                                                      [[error localizedDescription] UTF8String] :
+                                                      nullptr;
+                                error_str = err ? err : "nil";
+
+                                compilation_finished = true;
+                              }];
+
+      /* Immediately wait for either the compilation to finish or for app shutdown. */
+      while (ShaderCache::running && !compilation_finished) {
+        std::this_thread::sleep_for(std::chrono::milliseconds(5));
+      }
+    }
+
+    if (creating_new_archive && pipeline) {
+      /* Add pipeline into the new archive. */
+      NSError *error;
+      if (![archive addComputePipelineFunctionsWithDescriptor:computePipelineStateDescriptor
+                                                        error:&error])
+      {
+        NSString *errStr = [error localizedDescription];
+        metal_printf("Failed to add PSO to archive:\n%s", errStr ? [errStr UTF8String] : "nil");
+      }
+    }
+
+    if (!pipeline) {
+      metal_printf(
+          "newComputePipelineStateWithDescriptor failed for \"%s\"%s. "
+          "Error:\n%s\n",
+          device_kernel_as_string(device_kernel),
+          (archive && !recreate_archive) ? " Archive may be incomplete or corrupt - attempting "
+                                           "recreation.." :
+                                           "",
+          error_str.c_str());
+    }
+  };
+
+  double starttime = time_dt();
+
+  do_compilation();
+
+  /* An archive might have a corrupt entry and fail to materialize the pipeline. This shouldn't
+   * happen, but if it does we recreate it. */
+  if (pipeline == nil && archive) {
+    recreate_archive = true;
+    pipelineOptions = MTLPipelineOptionNone;
+    path_remove(metalbin_path);
+
+    do_compilation();
+  }
+
+  double duration = time_dt() - starttime;
+
+  if (pipeline == nil) {
+    metal_printf("%16s | %2d | %-55s | %7.2fs | FAILED!",
+                 kernel_type_as_string(pso_type),
+                 device_kernel,
+                 device_kernel_as_string(device_kernel),
+                 duration);
+    return;
+  }
+
+  if (!num_threads_per_block) {
+    num_threads_per_block = round_down(pipeline.maxTotalThreadsPerThreadgroup,
+                                       pipeline.threadExecutionWidth);
+    num_threads_per_block = std::max(num_threads_per_block, (int)pipeline.threadExecutionWidth);
+  }
+
+  if (ShaderCache::running) {
+    if (creating_new_archive || recreate_archive) {
+      if (![archive serializeToURL:[NSURL fileURLWithPath:@(metalbin_path.c_str())] error:&error])
+      {
+        metal_printf("Failed to save binary archive to %s, error:\n%s",
+                     metalbin_path.c_str(),
+                     [[error localizedDescription] UTF8String]);
+      }
+      else {
+        path_cache_kernel_mark_added_and_clear_old(metalbin_path);
+      }
+    }
+  }
+
+  this->loaded = true;
+  [computePipelineStateDescriptor release];
+  computePipelineStateDescriptor = nil;
+
+  if (!use_binary_archive) {
+    metal_printf("%16s | %2d | %-55s | %7.2fs",
+                 kernel_type_as_string(pso_type),
+                 int(device_kernel),
+                 device_kernel_as_string(device_kernel),
+                 duration);
+  }
+  else {
+    metal_printf("%16s | %2d | %-55s | %7.2fs | %s: %s",
+                 kernel_type_as_string(pso_type),
+                 device_kernel,
+                 device_kernel_as_string(device_kernel),
+                 duration,
+                 creating_new_archive ? " new" : "load",
+                 metalbin_name.c_str());
+  }
+}
+
+bool MetalDeviceKernels::load(MetalDevice *device, MetalPipelineType pso_type)
+{
+  auto *shader_cache = get_shader_cache(device->mtlDevice);
+  for (int i = 0; i < DEVICE_KERNEL_NUM; i++) {
+    shader_cache->load_kernel((DeviceKernel)i, device, pso_type);
+  }
+  return true;
+}
+
+void MetalDeviceKernels::wait_for_all()
+{
+  for (int i = 0; i < g_shaderCacheCount; i++) {
+    g_shaderCache[i].second->wait_for_all();
+  }
+}
+
+int MetalDeviceKernels::num_incomplete_specialization_requests()
+{
+  /* Return true if any ShaderCaches have ongoing specialization requests (typically there will be
+   * only 1). */
+  int total = 0;
+  for (int i = 0; i < g_shaderCacheCount; i++) {
+    total += g_shaderCache[i].second->incomplete_specialization_requests;
+  }
+  return total;
+}
+
+int MetalDeviceKernels::get_loaded_kernel_count(const MetalDevice *device,
+                                                MetalPipelineType pso_type)
+{
+  auto *shader_cache = get_shader_cache(device->mtlDevice);
+  int loaded_count = DEVICE_KERNEL_NUM;
+  for (int i = 0; i < DEVICE_KERNEL_NUM; i++) {
+    if (shader_cache->should_load_kernel((DeviceKernel)i, device, pso_type)) {
+      loaded_count -= 1;
+    }
+  }
+  return loaded_count;
+}
+
+bool MetalDeviceKernels::should_load_kernels(const MetalDevice *device, MetalPipelineType pso_type)
+{
+  return get_loaded_kernel_count(device, pso_type) != DEVICE_KERNEL_NUM;
+}
+
+const MetalKernelPipeline *MetalDeviceKernels::get_best_pipeline(const MetalDevice *device,
+                                                                 DeviceKernel kernel)
+{
+  return get_shader_cache(device->mtlDevice)->get_best_pipeline(kernel, device);
+}
+
+bool MetalDeviceKernels::is_benchmark_warmup()
+{
+  NSArray *args = [[NSProcessInfo processInfo] arguments];
+  for (int i = 0; i < args.count; i++) {
+    if (const char *arg = [[args objectAtIndex:i] cStringUsingEncoding:NSASCIIStringEncoding]) {
+      if (!strcmp(arg, "--warm-up")) {
+        return true;
+      }
+    }
+  }
+  return false;
+}
+
+void MetalDeviceKernels::static_deinitialize()
+{
+  for (int i = 0; i < g_shaderCacheCount; i++) {
+    g_shaderCache[i] = DeviceShaderCache();
+  }
+}
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/queue.h

@@ -0,0 +1,130 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#ifdef WITH_METAL
+
+#  include "device/kernel.h"
+#  include "device/memory.h"
+#  include "device/queue.h"
+
+#  include "device/metal/util.h"
+#  include "kernel/device/metal/globals.h"
+
+#  define MAX_SAMPLE_BUFFER_LENGTH 4096
+
+/* The number of resources to be contiguously encoded into the MetalAncillaries struct. */
+#  define ANCILLARY_SLOT_COUNT 11
+
+CCL_NAMESPACE_BEGIN
+
+class MetalDevice;
+
+/* Base class for Metal queues. */
+class MetalDeviceQueue : public DeviceQueue {
+ public:
+  MetalDeviceQueue(MetalDevice *device);
+  ~MetalDeviceQueue() override;
+
+  int num_concurrent_states(const size_t /*state_size*/) const override;
+  int num_concurrent_busy_states(const size_t /*state_size*/) const override;
+  int num_sort_partitions(int max_num_paths, uint max_scene_shaders) const override;
+  bool supports_local_atomic_sort() const override;
+
+  void init_execution() override;
+  void load_image_info() override;
+
+  bool enqueue(DeviceKernel kernel,
+               const int work_size,
+               const DeviceKernelArguments &args) override;
+
+  bool synchronize() override;
+
+  void zero_to_device(device_memory &mem) override;
+  void copy_to_device(device_memory &mem) override;
+  void copy_from_device(device_memory &mem) override;
+  void *copy_from_device_synchronized(device_memory &mem, vector<uint8_t> &storage) override;
+
+  void *native_queue() override;
+
+  unique_ptr<DeviceGraphicsInterop> graphics_interop_create() override;
+
+ protected:
+  void setup_capture();
+  void update_capture(DeviceKernel kernel);
+  void begin_capture();
+  void end_capture();
+  void prepare_resources();
+
+  id<MTLComputeCommandEncoder> get_compute_encoder(DeviceKernel kernel);
+  id<MTLBlitCommandEncoder> get_blit_encoder();
+
+  MetalDevice *metal_device_;
+
+  API_AVAILABLE(macos(11.0), ios(14.0))
+  MTLCommandBufferDescriptor *command_buffer_desc_ = nullptr;
+  id<MTLDevice> mtlDevice_ = nil;
+  id<MTLCommandQueue> mtlCommandQueue_ = nil;
+  id<MTLCommandBuffer> mtlCommandBuffer_ = nil;
+  id<MTLComputeCommandEncoder> mtlComputeEncoder_ = nil;
+  id<MTLBlitCommandEncoder> mtlBlitEncoder_ = nil;
+  API_AVAILABLE(macos(10.14), ios(14.0))
+  id<MTLSharedEvent> shared_event_ = nil;
+  API_AVAILABLE(macos(10.14), ios(14.0))
+  MTLSharedEventListener *shared_event_listener_ = nil;
+  MetalDispatchPipeline active_pipelines_[DEVICE_KERNEL_NUM];
+
+  dispatch_queue_t event_queue_;
+  dispatch_semaphore_t wait_semaphore_;
+
+  uint64_t shared_event_id_;
+  uint64_t command_buffers_submitted_ = 0;
+  uint64_t command_buffers_completed_ = 0;
+  Stats &stats_;
+
+  void close_compute_encoder();
+  void close_blit_encoder();
+
+  bool verbose_tracing_ = false;
+  bool label_command_encoders_ = false;
+
+  /* Per-kernel profiling (see CYCLES_METAL_PROFILING). */
+
+  struct TimingData {
+    DeviceKernel kernel;
+    int work_size;
+    uint64_t timing_id;
+  };
+  std::vector<TimingData> command_encoder_labels_;
+  bool profiling_enabled_ = false;
+  uint64_t current_encoder_idx_ = 0;
+
+  std::atomic<uint64_t> counter_sample_buffer_curr_idx_ = 0;
+
+  void flush_timing_stats();
+
+  struct TimingStats {
+    double total_time = 0.0;
+    uint64_t total_work_size = 0;
+    uint64_t num_dispatches = 0;
+  };
+  TimingStats timing_stats_[DEVICE_KERNEL_NUM];
+  double last_completion_time_ = 0.0;
+
+  /* .gputrace capture (see CYCLES_DEBUG_METAL_CAPTURE_...). */
+
+  id<MTLCaptureScope> mtlCaptureScope_ = nil;
+  DeviceKernel capture_kernel_;
+  int capture_dispatch_counter_ = 0;
+  bool capture_samples_ = false;
+  int capture_reset_counter_ = 0;
+  bool is_capturing_ = false;
+  bool is_capturing_to_disk_ = false;
+  bool has_captured_to_disk_ = false;
+};
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/queue.mm

@@ -0,0 +1,922 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#ifdef WITH_METAL
+
+#  include <algorithm>
+#  include <mutex>
+
+#  include "device/metal/queue.h"
+
+#  include "device/metal/device_impl.h"
+#  include "device/metal/graphics_interop.h"
+#  include "device/metal/kernel.h"
+
+#  include "util/path.h"
+#  include "util/string.h"
+#  include "util/time.h"
+
+CCL_NAMESPACE_BEGIN
+
+/* MetalDeviceQueue */
+
+MetalDeviceQueue::MetalDeviceQueue(MetalDevice *device)
+    : DeviceQueue(device), metal_device_(device), stats_(device->stats)
+{
+  @autoreleasepool {
+    command_buffer_desc_ = [[MTLCommandBufferDescriptor alloc] init];
+    command_buffer_desc_.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
+
+    mtlDevice_ = device->mtlDevice;
+    mtlCommandQueue_ = device->mtlComputeCommandQueue;
+
+    shared_event_ = [mtlDevice_ newSharedEvent];
+    shared_event_id_ = 1;
+
+    /* Shareable event listener */
+    event_queue_ = dispatch_queue_create("com.cycles.metal.event_queue", nullptr);
+    shared_event_listener_ = [[MTLSharedEventListener alloc] initWithDispatchQueue:event_queue_];
+
+    wait_semaphore_ = dispatch_semaphore_create(0);
+
+    if (auto *str = getenv("CYCLES_METAL_PROFILING")) {
+      if (atoi(str) && [mtlDevice_ supportsCounterSampling:MTLCounterSamplingPointAtStageBoundary])
+      {
+        /* Enable per-kernel timing breakdown (shown at end of render). */
+        profiling_enabled_ = true;
+        label_command_encoders_ = true;
+      }
+    }
+    if (getenv("CYCLES_METAL_DEBUG")) {
+      /* Enable very verbose tracing (shows every dispatch). */
+      verbose_tracing_ = true;
+      label_command_encoders_ = true;
+    }
+
+    setup_capture();
+  }
+}
+
+void MetalDeviceQueue::setup_capture()
+{
+  capture_kernel_ = DeviceKernel(-1);
+
+  if (auto *capture_kernel_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_KERNEL")) {
+    /* CYCLES_DEBUG_METAL_CAPTURE_KERNEL captures a single dispatch of the specified kernel. */
+    capture_kernel_ = DeviceKernel(atoi(capture_kernel_str));
+    printf("Capture kernel: %d = %s\n", capture_kernel_, device_kernel_as_string(capture_kernel_));
+
+    capture_dispatch_counter_ = 0;
+    if (auto *capture_dispatch_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_DISPATCH")) {
+      capture_dispatch_counter_ = atoi(capture_dispatch_str);
+
+      printf("Capture dispatch number %d\n", capture_dispatch_counter_);
+    }
+  }
+  else if (auto *capture_samples_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_SAMPLES")) {
+    /* CYCLES_DEBUG_METAL_CAPTURE_SAMPLES captures a block of dispatches from reset#(N) to
+     * reset#(N+1). */
+    capture_samples_ = true;
+    capture_reset_counter_ = atoi(capture_samples_str);
+
+    capture_dispatch_counter_ = INT_MAX;
+    if (auto *capture_limit_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_LIMIT")) {
+      /* CYCLES_DEBUG_METAL_CAPTURE_LIMIT sets the maximum number of dispatches to capture. */
+      capture_dispatch_counter_ = atoi(capture_limit_str);
+    }
+
+    printf("Capturing sample block %d (dispatch limit: %d)\n",
+           capture_reset_counter_,
+           capture_dispatch_counter_);
+  }
+  else {
+    /* No capturing requested. */
+    return;
+  }
+
+  /* Enable .gputrace capture for the specified DeviceKernel. */
+  MTLCaptureManager *captureManager = [MTLCaptureManager sharedCaptureManager];
+  mtlCaptureScope_ = [captureManager newCaptureScopeWithDevice:mtlDevice_];
+  mtlCaptureScope_.label = [NSString stringWithFormat:@"Cycles kernel dispatch"];
+  [captureManager setDefaultCaptureScope:mtlCaptureScope_];
+
+  label_command_encoders_ = true;
+
+  if (auto *capture_url = getenv("CYCLES_DEBUG_METAL_CAPTURE_URL")) {
+    if ([captureManager supportsDestination:MTLCaptureDestinationGPUTraceDocument]) {
+
+      MTLCaptureDescriptor *captureDescriptor = [[MTLCaptureDescriptor alloc] init];
+      captureDescriptor.captureObject = mtlCaptureScope_;
+      captureDescriptor.destination = MTLCaptureDestinationGPUTraceDocument;
+      captureDescriptor.outputURL = [NSURL fileURLWithPath:@(capture_url)];
+
+      NSError *error;
+      if (![captureManager startCaptureWithDescriptor:captureDescriptor error:&error]) {
+        NSString *err = [error localizedDescription];
+        printf("Start capture failed: %s\n", [err UTF8String]);
+      }
+      else {
+        printf("Capture started (URL: %s)\n", capture_url);
+        is_capturing_to_disk_ = true;
+      }
+    }
+    else {
+      printf("Capture to file is not supported\n");
+    }
+  }
+}
+
+void MetalDeviceQueue::update_capture(DeviceKernel kernel)
+{
+  /* Handle capture end triggers. */
+  if (is_capturing_) {
+    capture_dispatch_counter_ -= 1;
+    if (capture_dispatch_counter_ <= 0 || kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
+      /* End capture if we've hit the dispatch limit or we hit a "reset". */
+      end_capture();
+    }
+    return;
+  }
+
+  if (capture_dispatch_counter_ < 0) {
+    /* We finished capturing. */
+    return;
+  }
+
+  /* Handle single-capture start trigger. */
+  if (kernel == capture_kernel_) {
+    /* Start capturing when we hit the Nth dispatch of the specified kernel. */
+    if (capture_dispatch_counter_ == 0) {
+      begin_capture();
+    }
+    capture_dispatch_counter_ -= 1;
+    return;
+  }
+
+  /* Handle multi-capture start trigger. */
+  if (capture_samples_) {
+    /* Start capturing when the reset countdown is at 0. */
+    if (capture_reset_counter_ == 0) {
+      begin_capture();
+    }
+
+    if (kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
+      capture_reset_counter_ -= 1;
+    }
+    return;
+  }
+}
+
+void MetalDeviceQueue::begin_capture()
+{
+  /* Start gputrace capture. */
+  if (mtlCommandBuffer_) {
+    synchronize();
+  }
+  [mtlCaptureScope_ beginScope];
+  printf("[mtlCaptureScope_ beginScope]\n");
+  is_capturing_ = true;
+}
+
+void MetalDeviceQueue::end_capture()
+{
+  [mtlCaptureScope_ endScope];
+  is_capturing_ = false;
+  printf("[mtlCaptureScope_ endScope]\n");
+
+  if (is_capturing_to_disk_) {
+    [[MTLCaptureManager sharedCaptureManager] stopCapture];
+    has_captured_to_disk_ = true;
+    is_capturing_to_disk_ = false;
+    is_capturing_ = false;
+    printf("Capture stopped\n");
+  }
+}
+
+MetalDeviceQueue::~MetalDeviceQueue()
+{
+  /* Tidying up here isn't really practical - we should expect and require the work
+   * queue to be empty here. */
+  assert(mtlCommandBuffer_ == nil);
+  assert(command_buffers_submitted_ == command_buffers_completed_);
+
+  close_compute_encoder();
+  close_blit_encoder();
+
+  [shared_event_listener_ release];
+  [shared_event_ release];
+  [command_buffer_desc_ release];
+
+  if (mtlCaptureScope_) {
+    [mtlCaptureScope_ release];
+  }
+
+  double total_time = 0.0;
+
+  /* Show per-kernel timings, if gathered (see CYCLES_METAL_PROFILING). */
+  int64_t num_dispatches = 0;
+  int64_t num_pathtracing_dispatches = 0;
+  for (size_t i = 0; i < DEVICE_KERNEL_NUM; i++) {
+    auto &stat = timing_stats_[i];
+    bool pathtracing_kernel = (i <= DEVICE_KERNEL_INTEGRATOR_RESET) &&
+                              (i != DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL);
+    total_time += stat.total_time;
+    num_dispatches += stat.num_dispatches;
+    num_pathtracing_dispatches += pathtracing_kernel ? stat.num_dispatches : 0;
+  }
+  bool has_extra = (num_pathtracing_dispatches && num_dispatches > num_pathtracing_dispatches);
+
+  if (num_dispatches) {
+    printf("\nMetal %sdispatch stats:\n", num_pathtracing_dispatches ? "path-tracing " : "");
+    auto header = string_printf("%-40s %16s %12s %12s %9s %9s",
+                                "Kernel name",
+                                "Total threads",
+                                "Dispatches",
+                                "Avg. T/D",
+                                "Time/s",
+                                "Time/%");
+    auto divider = string(header.length(), '-');
+    printf("%s\n%s\n%s\n", divider.c_str(), header.c_str(), divider.c_str());
+
+    for (size_t i = 0; i < DEVICE_KERNEL_NUM; i++) {
+      auto &stat = timing_stats_[i];
+
+      bool pathtracing_kernel = (i <= DEVICE_KERNEL_INTEGRATOR_RESET) &&
+                                (i != DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL);
+      if ((pathtracing_kernel && num_pathtracing_dispatches) || stat.num_dispatches > 0) {
+        printf("%-40s %16llu %12llu %12llu %9.4f %9.2f\n",
+               device_kernel_as_string(DeviceKernel(i)),
+               stat.total_work_size,
+               stat.num_dispatches,
+               stat.total_work_size / stat.num_dispatches,
+               stat.total_time,
+               stat.total_time * 100.0 / total_time);
+      }
+      if (has_extra && i == DEVICE_KERNEL_INTEGRATOR_RESET) {
+        printf("%s\n", divider.c_str());
+      }
+    }
+    printf("%s\n", divider.c_str());
+    printf("%-40s %16s %12llu %12s %9.4f %9.2f\n", "", "", num_dispatches, "", total_time, 100.0);
+    printf("%s\n\n", divider.c_str());
+  }
+}
+
+int MetalDeviceQueue::num_concurrent_states(const size_t state_size) const
+{
+  size_t state_count = 4194304;
+
+  /* Increasing the state count doesn't notably benefit M1-family systems. */
+  if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) != APPLE_M1) {
+    const size_t max_recommended_working_set =
+        [metal_device_->mtlDevice recommendedMaxWorkingSetSize];
+
+    /* Only use 90% of available working set for safety. */
+    size_t percent = 90;
+    if (auto *str = getenv("CYCLES_METAL_WORKING_SET_PERCENT")) {
+      percent = atoi(str);
+    }
+
+    const size_t max_working_set = (max_recommended_working_set * percent) / 100;
+    size_t max_safe_state_count = 0;
+
+    if (stats_.mem_used < max_working_set) {
+      const size_t headroom = max_working_set - stats_.mem_used;
+      max_safe_state_count = headroom / state_size;
+    }
+
+    /* Require a bare minimum of states to avoid pathological performance. */
+    if (max_safe_state_count >= 65536) {
+      /* If RAM is limited, we can still render with reduced state count. */
+      if (max_safe_state_count < state_count) {
+        metal_printf(
+            "Reducing state count to fit within available RAM. %zu -> %zu (%.1f%% of original "
+            "size)",
+            state_count,
+            max_safe_state_count,
+            double(max_safe_state_count) / double(state_count) * 100.0);
+        state_count = max_safe_state_count;
+      }
+      else {
+        /* Aggressive safety margin: only grow if it leaves us at < 50% max working set
+         * utilization. */
+        size_t grow_percent = 50;
+        if (auto *str = getenv("CYCLES_METAL_GROW_PERCENT")) {
+          grow_percent = atoi(str);
+        }
+
+        max_safe_state_count = (max_safe_state_count * grow_percent) / 100;
+
+        /* Limit to two "doublings" - we see diminishing returns after that. */
+        for (int i = 0; i < 2; i++) {
+          /* Determine whether we can double the state count, and leave enough GPU-available
+           * memory. Enlarging the state size allows us to keep dispatch sizes high and minimize
+           * work submission overheads. */
+          if (max_safe_state_count > state_count * 2) {
+            state_count *= 2;
+            metal_printf("Doubling state count to exploit available RAM (new size = %zu)",
+                         state_count);
+          }
+        }
+      }
+    }
+    else {
+      metal_device_->set_error("Out of memory - couldn't allocate integrator state");
+      state_count = 0;
+    }
+  }
+  return state_count;
+}
+
+int MetalDeviceQueue::num_concurrent_busy_states(const size_t state_size) const
+{
+  /* A 1:4 busy:total ratio gives best rendering performance, independent of total state count. */
+  return num_concurrent_states(state_size) / 4;
+}
+
+int MetalDeviceQueue::num_sort_partitions(int max_num_paths, uint max_scene_shaders) const
+{
+  int sort_partition_elements = MetalInfo::optimal_sort_partition_elements();
+  /* Sort partitioning becomes less effective when more shaders are in the wavefront. In lieu of
+   * a more sophisticated heuristic we simply disable sort partitioning if the shader count is
+   * high.
+   */
+  if (max_scene_shaders < 300 && sort_partition_elements > 0) {
+    return max(max_num_paths / sort_partition_elements, 1);
+  }
+  else {
+    return 1;
+  }
+}
+
+bool MetalDeviceQueue::supports_local_atomic_sort() const
+{
+  return metal_device_->use_local_atomic_sort();
+}
+
+static void zero_resource(void *address_in_arg_buffer, int index = 0)
+{
+  uint64_t *pptr = (uint64_t *)address_in_arg_buffer;
+  pptr[index] = 0;
+}
+
+template<class T> void write_resource(void *address_in_arg_buffer, T resource, int index = 0)
+{
+  zero_resource(address_in_arg_buffer, index);
+  uint64_t *pptr = (uint64_t *)address_in_arg_buffer;
+  if (resource) {
+    pptr[index] = metal_gpuResourceID(resource);
+  }
+}
+
+template<> void write_resource(void *address_in_arg_buffer, id<MTLBuffer> buffer, int index)
+{
+  zero_resource(address_in_arg_buffer, index);
+  uint64_t *pptr = (uint64_t *)address_in_arg_buffer;
+  if (buffer) {
+    pptr[index] = metal_gpuAddress(buffer);
+  }
+}
+
+static id<MTLBuffer> patch_resource(void *address_in_arg_buffer, int index = 0)
+{
+  uint64_t *pptr = (uint64_t *)address_in_arg_buffer;
+  if (MetalDevice::MetalMem *mmem = (MetalDevice::MetalMem *)pptr[index]) {
+    write_resource<id<MTLBuffer>>(address_in_arg_buffer, mmem->mtlBuffer, index);
+    return mmem->mtlBuffer;
+  }
+  return nil;
+}
+
+void MetalDeviceQueue::init_execution()
+{
+  /* Populate blas_array. */
+  uint64_t *blas_array = (uint64_t *)metal_device_->blas_buffer.contents;
+  for (uint64_t slot = 0; slot < metal_device_->blas_array.size(); ++slot) {
+    write_resource(blas_array, metal_device_->blas_array[slot], slot);
+  }
+
+  /* Populate image bindings. */
+  load_image_info();
+
+  /* Synchronize memory copies. */
+  synchronize();
+}
+
+void MetalDeviceQueue::load_image_info()
+{
+  /* TODO: Can this be optimized to only update info ids that changed? Why is this done delayed
+   * instead of immediately when allocating the image? */
+  device_vector<KernelImageInfo> &image_info = metal_device_->image_info;
+  id<MTLBuffer> &image_bindings = metal_device_->image_bindings;
+  std::vector<id<MTLResource>> &image_info_id_map = metal_device_->image_info_id_map;
+
+  /* Ensure image_info is allocated before populating. */
+  image_info.copy_to_device();
+
+  /* Populate texture bindings. */
+  uint64_t *bindings = (uint64_t *)image_bindings.contents;
+  memset(bindings, 0, image_bindings.length);
+  for (int image_info_id = 0; image_info_id < image_info.size(); ++image_info_id) {
+    if (image_info_id_map[image_info_id]) {
+      if (metal_device_->is_texture(image_info[image_info_id])) {
+        write_resource(bindings, id<MTLTexture>(image_info_id_map[image_info_id]), image_info_id);
+      }
+      else {
+        /* The GPU address of a 1D buffer texture is written into the image_info_id data field. */
+        write_resource(
+            &image_info[image_info_id].data, id<MTLBuffer>(image_info_id_map[image_info_id]), 0);
+      }
+    }
+  }
+}
+
+bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
+                               const int work_size,
+                               const DeviceKernelArguments &args)
+{
+  @autoreleasepool {
+    update_capture(kernel);
+
+    if (metal_device_->have_error()) {
+      return false;
+    }
+
+    debug_enqueue_begin(kernel, work_size);
+
+    LOG_TRACE << "Metal queue launch " << device_kernel_as_string(kernel) << ", work_size "
+              << work_size;
+
+    id<MTLComputeCommandEncoder> mtlComputeCommandEncoder = get_compute_encoder(kernel);
+
+    if (profiling_enabled_) {
+      command_encoder_labels_.push_back({kernel, work_size, current_encoder_idx_});
+    }
+    if (label_command_encoders_) {
+      /* Add human-readable labels if we're doing any form of debugging / profiling. */
+      mtlComputeCommandEncoder.label = [NSString
+          stringWithFormat:@"Metal queue launch %s, work_size %d",
+                           device_kernel_as_string(kernel),
+                           work_size];
+    }
+
+    if (!active_pipelines_[kernel].update(metal_device_, kernel)) {
+      metal_device_->set_error(
+          string_printf("Could not activate pipeline for %s\n", device_kernel_as_string(kernel)));
+      return false;
+    }
+    MetalDispatchPipeline &active_pipeline = active_pipelines_[kernel];
+
+    uint8_t dynamic_args[512] = {0};
+
+    /* Prepare the dynamic "enqueue" arguments */
+    size_t dynamic_bytes_written = 0;
+    size_t max_size_in_bytes = 0;
+    for (size_t i = 0; i < args.count; i++) {
+      size_t size_in_bytes = args.sizes[i];
+      max_size_in_bytes = max(max_size_in_bytes, size_in_bytes);
+      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) {
+        patch_resource(dynamic_args + dynamic_bytes_written);
+      }
+      dynamic_bytes_written += size_in_bytes;
+    }
+    /* Apply conventional struct alignment (stops asserts firing when API validation is enabled).
+     */
+    dynamic_bytes_written = round_up(dynamic_bytes_written, max_size_in_bytes);
+
+    /* Check that the dynamic args didn't overflow. */
+    assert(dynamic_bytes_written <= sizeof(dynamic_args));
+
+    uint64_t ancillary_args[ANCILLARY_SLOT_COUNT] = {0};
+
+    /* Encode ancillaries */
+    int ancillary_index = 0;
+    write_resource(ancillary_args, metal_device_->image_bindings, ancillary_index++);
+
+    if (metal_device_->use_metalrt) {
+      write_resource(ancillary_args, metal_device_->accel_struct, ancillary_index++);
+      write_resource(ancillary_args, metal_device_->blas_buffer, ancillary_index++);
+
+      /* Write the intersection function table. */
+      for (int table_idx = 0; table_idx < METALRT_TABLE_NUM; table_idx++) {
+        write_resource(
+            ancillary_args, active_pipeline.intersection_func_table[table_idx], ancillary_index++);
+      }
+      assert(ancillary_index == ANCILLARY_SLOT_COUNT);
+    }
+
+    [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];
+
+    /* 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) {
+          /* Mark all Accelerations resources as used */
+          [mtlComputeCommandEncoder useResource:accel_struct usage:MTLResourceUsageRead];
+          if (metal_device_->blas_buffer) {
+            [mtlComputeCommandEncoder useResource:metal_device_->blas_buffer
+                                            usage:MTLResourceUsageRead];
+          }
+          [mtlComputeCommandEncoder useResources:metal_device_->unique_blas_array.data()
+                                           count:metal_device_->unique_blas_array.size()
+                                           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];
+
+    /* Compute kernel launch parameters. */
+    const int num_threads_per_block = active_pipeline.num_threads_per_block;
+
+    int shared_mem_bytes = 0;
+
+    switch (kernel) {
+      case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_QUEUED_SHADOW_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_ACTIVE_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY:
+      case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY:
+        /* See parallel_active_index.h for why this amount of shared memory is needed.
+         * Rounded up to 16 bytes for Metal */
+        shared_mem_bytes = (int)round_up((num_threads_per_block + 1) * sizeof(int), 16);
+        break;
+
+      case DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS:
+      case DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS: {
+        int key_count = metal_device_->launch_params->data.max_shaders;
+        shared_mem_bytes = (int)round_up(key_count * sizeof(int), 16);
+        break;
+      }
+
+      default:
+        break;
+    }
+
+    if (shared_mem_bytes) {
+      assert(shared_mem_bytes <= 32 * 1024);
+      [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
+    }
+
+    MTLSize size_threads_per_dispatch = MTLSizeMake(work_size, 1, 1);
+    MTLSize size_threads_per_threadgroup = MTLSizeMake(num_threads_per_block, 1, 1);
+    [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;
+      if (command_buffer.status != MTLCommandBufferStatusCompleted) {
+        str = string_printf("Command buffer not completed. status = %d. ",
+                            int(command_buffer.status));
+      }
+      if (command_buffer.error) {
+        @autoreleasepool {
+          const char *errCStr = [[NSString stringWithFormat:@"%@", command_buffer.error]
+              UTF8String];
+          str += string_printf("(%s.%s):\n%s\n",
+                               kernel_type_as_string(active_pipeline.pso_type),
+                               device_kernel_as_string(kernel),
+                               errCStr);
+        }
+      }
+      if (!str.empty()) {
+        metal_device_->set_error(str);
+      }
+    }];
+
+    if (verbose_tracing_ || is_capturing_) {
+      /* Force a sync we've enabled step-by-step verbose tracing or if we're capturing. */
+      synchronize();
+
+      /* Show queue counters and dispatch timing. */
+      if (verbose_tracing_) {
+        if (kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
+          printf(
+              "_____________________________________.____________________.______________._________"
+              "__"
+              "______________________________________\n");
+        }
+
+        printf("%-40s| %7d threads |%5.2fms | buckets [",
+               device_kernel_as_string(kernel),
+               work_size,
+               last_completion_time_ * 1000.0);
+        std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+        for (auto &it : metal_device_->metal_mem_map) {
+          const string c_integrator_queue_counter = "integrator_queue_counter";
+          if (it.first->global_name() == c_integrator_queue_counter) {
+            if (IntegratorQueueCounter *queue_counter = (IntegratorQueueCounter *)
+                                                            it.first->host_pointer)
+            {
+              for (int i = 0; i < DEVICE_GPU_KERNEL_INTEGRATOR_NUM; i++) {
+                printf("%s%d", i == 0 ? "" : ",", queue_counter->num_queued[i]);
+              }
+            }
+            break;
+          }
+        }
+        printf("]\n");
+      }
+    }
+
+    debug_enqueue_end();
+
+    return !(metal_device_->have_error());
+  }
+}
+
+void MetalDeviceQueue::flush_timing_stats()
+{
+  for (auto label : command_encoder_labels_) {
+    TimingStats &stat = timing_stats_[label.kernel];
+
+    double completion_time_gpu;
+    NSData *computeTimeStamps = [metal_device_->mtlCounterSampleBuffer
+        resolveCounterRange:NSMakeRange(label.timing_id, 2)];
+    MTLCounterResultTimestamp *timestamps = (MTLCounterResultTimestamp *)(computeTimeStamps.bytes);
+
+    uint64_t begTime = timestamps[0].timestamp;
+    uint64_t endTime = timestamps[1].timestamp;
+    completion_time_gpu = (endTime - begTime) / (double)NSEC_PER_SEC;
+
+    stat.num_dispatches++;
+    stat.total_time += completion_time_gpu;
+    stat.total_work_size += label.work_size;
+    last_completion_time_ = completion_time_gpu;
+  }
+  command_encoder_labels_.clear();
+}
+
+bool MetalDeviceQueue::synchronize()
+{
+  @autoreleasepool {
+    if (has_captured_to_disk_ || metal_device_->have_error()) {
+      return false;
+    }
+
+    close_compute_encoder();
+    close_blit_encoder();
+
+    if (mtlCommandBuffer_) {
+      scoped_timer timer;
+
+      uint64_t shared_event_id_ = this->shared_event_id_++;
+
+      __block dispatch_semaphore_t block_sema = wait_semaphore_;
+      [shared_event_ notifyListener:shared_event_listener_
+                            atValue:shared_event_id_
+                              block:^(id<MTLSharedEvent> /*sharedEvent*/, uint64_t /*value*/) {
+                                dispatch_semaphore_signal(block_sema);
+                              }];
+
+      [mtlCommandBuffer_ encodeSignalEvent:shared_event_ value:shared_event_id_];
+      [mtlCommandBuffer_ commit];
+      dispatch_semaphore_wait(wait_semaphore_, DISPATCH_TIME_FOREVER);
+
+      [mtlCommandBuffer_ release];
+
+      metal_device_->flush_delayed_free_list();
+
+      mtlCommandBuffer_ = nil;
+      flush_timing_stats();
+    }
+
+    debug_synchronize();
+
+    return !(metal_device_->have_error());
+  }
+}
+
+void MetalDeviceQueue::zero_to_device(device_memory &mem)
+{
+  @autoreleasepool {
+    if (metal_device_->have_error()) {
+      return;
+    }
+
+    assert(mem.type != MEM_IMAGE_TEXTURE);
+
+    if (mem.memory_size() == 0) {
+      return;
+    }
+
+    /* Allocate on demand. */
+    if (mem.device_pointer == 0) {
+      metal_device_->mem_alloc(mem);
+    }
+
+    /* Zero memory on device. */
+    assert(mem.device_pointer != 0);
+
+    std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+    MetalDevice::MetalMem &mmem = *metal_device_->metal_mem_map.at(&mem);
+    if (mmem.mtlBuffer) {
+      id<MTLBlitCommandEncoder> blitEncoder = get_blit_encoder();
+      [blitEncoder fillBuffer:mmem.mtlBuffer range:NSMakeRange(mmem.offset, mmem.size) value:0];
+    }
+    else {
+      metal_device_->mem_zero(mem);
+    }
+  }
+}
+
+void MetalDeviceQueue::copy_to_device(device_memory &mem)
+{
+  @autoreleasepool {
+    if (metal_device_->have_error()) {
+      return;
+    }
+
+    if (mem.memory_size() == 0) {
+      return;
+    }
+
+    /* Allocate on demand. */
+    if (mem.device_pointer == 0) {
+      metal_device_->mem_alloc(mem);
+    }
+
+    assert(mem.device->mem_device_ptr(mem, metal_device_) != 0);
+    assert(mem.host_pointer != nullptr);
+    /* No need to copy - Apple Silicon has Unified Memory Architecture. */
+  }
+}
+
+void MetalDeviceQueue::copy_from_device(device_memory & /*mem*/)
+{
+  /* No need to copy - Apple Silicon has Unified Memory Architecture. */
+}
+
+void *MetalDeviceQueue::copy_from_device_synchronized(device_memory &mem,
+                                                      vector<uint8_t> & /*storage*/)
+{
+  if (mem.memory_size() == 0) {
+    return nullptr;
+  }
+
+  /* Wait until kernels have finished before returning from unified memory. */
+  synchronize();
+
+  device_ptr d_ptr = mem.device->mem_device_ptr(mem, metal_device_);
+  return (d_ptr) ? reinterpret_cast<MetalDevice::MetalMem *>(d_ptr)->hostPtr : nullptr;
+}
+
+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 */
+  for (auto &it : metal_device_->metal_mem_map) {
+    device_memory *mem = it.first;
+
+    MTLResourceUsage usage = MTLResourceUsageRead;
+    if (mem->type != MEM_GLOBAL && mem->type != MEM_READ_ONLY && mem->type != MEM_IMAGE_TEXTURE) {
+      usage |= MTLResourceUsageWrite;
+    }
+
+    if (it.second->mtlBuffer) {
+      /* METAL_WIP - use array version (i.e. useResources) */
+      [mtlComputeEncoder_ useResource:it.second->mtlBuffer usage:usage];
+    }
+    else if (it.second->mtlTexture) {
+      /* METAL_WIP - use array version (i.e. useResources) */
+      [mtlComputeEncoder_ useResource:it.second->mtlTexture usage:usage | MTLResourceUsageSample];
+    }
+  }
+
+  /* ancillaries */
+  [mtlComputeEncoder_ useResource:metal_device_->image_bindings usage:MTLResourceUsageRead];
+}
+
+id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel kernel)
+{
+  bool concurrent = int(kernel) < int(DEVICE_GPU_KERNEL_INTEGRATOR_NUM);
+
+  if (profiling_enabled_) {
+    /* Close the current encoder to ensure we're able to capture per-encoder timing data. */
+    close_compute_encoder();
+  }
+
+  if (mtlComputeEncoder_) {
+    if (mtlComputeEncoder_.dispatchType == concurrent ? MTLDispatchTypeConcurrent :
+                                                        MTLDispatchTypeSerial)
+    {
+      /* declare usage of MTLBuffers etc */
+      prepare_resources();
+
+      return mtlComputeEncoder_;
+    }
+    close_compute_encoder();
+  }
+
+  close_blit_encoder();
+
+  if (!mtlCommandBuffer_) {
+    mtlCommandBuffer_ = [mtlCommandQueue_ commandBuffer];
+    [mtlCommandBuffer_ retain];
+  }
+
+  if (profiling_enabled_) {
+    MTLComputePassDescriptor *desc = [[MTLComputePassDescriptor alloc] init];
+
+    current_encoder_idx_ = (counter_sample_buffer_curr_idx_.fetch_add(2) %
+                            MAX_SAMPLE_BUFFER_LENGTH);
+    [desc.sampleBufferAttachments[0] setSampleBuffer:metal_device_->mtlCounterSampleBuffer];
+    [desc.sampleBufferAttachments[0] setStartOfEncoderSampleIndex:current_encoder_idx_];
+    [desc.sampleBufferAttachments[0] setEndOfEncoderSampleIndex:current_encoder_idx_ + 1];
+
+    [desc setDispatchType:concurrent ? MTLDispatchTypeConcurrent : MTLDispatchTypeSerial];
+
+    mtlComputeEncoder_ = [mtlCommandBuffer_ computeCommandEncoderWithDescriptor:desc];
+  }
+  else {
+    mtlComputeEncoder_ = [mtlCommandBuffer_
+        computeCommandEncoderWithDispatchType:concurrent ? MTLDispatchTypeConcurrent :
+                                                           MTLDispatchTypeSerial];
+  }
+
+  [mtlComputeEncoder_ retain];
+  [mtlComputeEncoder_ setLabel:@(device_kernel_as_string(kernel))];
+
+  /* declare usage of MTLBuffers etc */
+  prepare_resources();
+
+  return mtlComputeEncoder_;
+}
+
+id<MTLBlitCommandEncoder> MetalDeviceQueue::get_blit_encoder()
+{
+  if (mtlBlitEncoder_) {
+    return mtlBlitEncoder_;
+  }
+
+  close_compute_encoder();
+
+  if (!mtlCommandBuffer_) {
+    mtlCommandBuffer_ = [mtlCommandQueue_ commandBuffer];
+    [mtlCommandBuffer_ retain];
+  }
+
+  mtlBlitEncoder_ = [mtlCommandBuffer_ blitCommandEncoder];
+  [mtlBlitEncoder_ retain];
+  return mtlBlitEncoder_;
+}
+
+void MetalDeviceQueue::close_compute_encoder()
+{
+  if (mtlComputeEncoder_) {
+    [mtlComputeEncoder_ endEncoding];
+    [mtlComputeEncoder_ release];
+    mtlComputeEncoder_ = nil;
+  }
+}
+
+void MetalDeviceQueue::close_blit_encoder()
+{
+  if (mtlBlitEncoder_) {
+    [mtlBlitEncoder_ endEncoding];
+    [mtlBlitEncoder_ release];
+    mtlBlitEncoder_ = nil;
+  }
+}
+
+void *MetalDeviceQueue::native_queue()
+{
+  return mtlCommandQueue_;
+}
+
+unique_ptr<DeviceGraphicsInterop> MetalDeviceQueue::graphics_interop_create()
+{
+  return make_unique<MetalDeviceGraphicsInterop>(this);
+}
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/util.h

@@ -0,0 +1,50 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#pragma once
+
+#ifdef WITH_METAL
+
+#  include <Metal/Metal.h>
+#  include <string>
+
+#  include "device/metal/device.h"
+#  include "device/metal/kernel.h"
+#  include "device/queue.h"
+
+#  include "util/thread.h"
+
+#  define metal_printf LOG_TRACE << string_printf
+
+CCL_NAMESPACE_BEGIN
+
+enum AppleGPUArchitecture {
+  APPLE_M1,
+  APPLE_M2,
+  APPLE_M2_BIG,
+  APPLE_M3,
+  /* Keep APPLE_UNKNOWN at the end of this enum to ensure that unknown future architectures get
+   * the most recent defaults when using comparison operators. */
+  APPLE_UNKNOWN,
+};
+
+/* Contains static Metal helper functions. */
+struct MetalInfo {
+  static const vector<id<MTLDevice>> &get_usable_devices();
+  static int get_apple_gpu_core_count(id<MTLDevice> device);
+  static AppleGPUArchitecture get_apple_gpu_architecture(id<MTLDevice> device);
+  static int optimal_sort_partition_elements();
+  static string get_device_name(id<MTLDevice> device);
+};
+
+void metal_gpu_address_helper_init(id<MTLDevice> device);
+
+uint64_t metal_gpuAddress(id<MTLBuffer> buffer);
+uint64_t metal_gpuResourceID(id<MTLTexture> texture);
+uint64_t metal_gpuResourceID(id<MTLAccelerationStructure> accel_struct);
+uint64_t metal_gpuResourceID(id<MTLIntersectionFunctionTable> ift);
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */

blender-5.2.0/intern/cycles/device/metal/util.mm

@@ -0,0 +1,232 @@
+/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
+ *
+ * SPDX-License-Identifier: Apache-2.0 */
+
+#ifdef WITH_METAL
+
+#  include "device/metal/util.h"
+#  include "device/metal/device_impl.h"
+#  include "util/md5.h"
+#  include "util/path.h"
+#  include "util/string.h"
+#  include "util/time.h"
+
+#  include <IOKit/IOKitLib.h>
+#  include <ctime>
+#  include <pwd.h>
+#  include <sys/shm.h>
+
+CCL_NAMESPACE_BEGIN
+
+/* Comment this out to test workaround for getting gpuAddress and gpuResourceID on macOS < 13.0. */
+#  define CYCLES_USE_TIER2D_BINDLESS
+
+string MetalInfo::get_device_name(id<MTLDevice> device)
+{
+  string device_name = [device.name UTF8String];
+
+  /* Append the GPU core count so we can distinguish between GPU variants in benchmarks. */
+  int gpu_core_count = get_apple_gpu_core_count(device);
+  device_name += string_printf(gpu_core_count ? " (GPU - %d cores)" : " (GPU)", gpu_core_count);
+
+  return device_name;
+}
+
+int MetalInfo::get_apple_gpu_core_count(id<MTLDevice> device)
+{
+  int core_count = 0;
+  if (@available(macos 12.0, *)) {
+    io_service_t gpu_service = IOServiceGetMatchingService(
+        kIOMainPortDefault, IORegistryEntryIDMatching(device.registryID));
+    if (CFNumberRef numberRef = (CFNumberRef)IORegistryEntryCreateCFProperty(
+            gpu_service, CFSTR("gpu-core-count"), nullptr, 0))
+    {
+      if (CFGetTypeID(numberRef) == CFNumberGetTypeID()) {
+        CFNumberGetValue(numberRef, kCFNumberSInt32Type, &core_count);
+      }
+      CFRelease(numberRef);
+    }
+  }
+  return core_count;
+}
+
+AppleGPUArchitecture MetalInfo::get_apple_gpu_architecture(id<MTLDevice> device)
+{
+  const char *device_name = [device.name UTF8String];
+  if (strstr(device_name, "M1")) {
+    return APPLE_M1;
+  }
+  if (strstr(device_name, "M2")) {
+    return get_apple_gpu_core_count(device) <= 10 ? APPLE_M2 : APPLE_M2_BIG;
+  }
+  if (strstr(device_name, "M3")) {
+    return APPLE_M3;
+  }
+  return APPLE_UNKNOWN;
+}
+
+int MetalInfo::optimal_sort_partition_elements()
+{
+  if (auto *str = getenv("CYCLES_METAL_SORT_PARTITION_ELEMENTS")) {
+    return atoi(str);
+  }
+
+  /* On M1 and M2 GPUs, we see better cache utilization if we partition the active indices before
+   * sorting each partition by material. Partitioning into chunks of 65536 elements results in an
+   * overall render time speedup of up to 15%. */
+
+  return 65536;
+}
+
+const vector<id<MTLDevice>> &MetalInfo::get_usable_devices()
+{
+  static vector<id<MTLDevice>> usable_devices;
+  static bool already_enumerated = false;
+
+  if (already_enumerated) {
+    return usable_devices;
+  }
+
+  metal_printf("Usable Metal devices:");
+  for (id<MTLDevice> device in MTLCopyAllDevices()) {
+    string device_name = get_device_name(device);
+    bool usable = false;
+
+    if (@available(macos 12.2, *)) {
+      const char *device_name_char = [device.name UTF8String];
+      if (!(strstr(device_name_char, "Intel") || strstr(device_name_char, "AMD")) &&
+          strstr(device_name_char, "Apple"))
+      {
+        /* TODO: Implement a better way to identify device vendor instead of relying on name. */
+        /* We only support Apple Silicon GPUs which all have unified memory, but explicitly check
+         * just in case it ever changes. */
+        usable = [device hasUnifiedMemory];
+      }
+    }
+
+    if (usable) {
+      metal_printf("- %s", device_name.c_str());
+      [device retain];
+      usable_devices.push_back(device);
+    }
+    else {
+      metal_printf("  (skipping \"%s\")", device_name.c_str());
+    }
+  }
+  if (usable_devices.empty()) {
+    metal_printf("   No usable Metal devices found");
+  }
+  already_enumerated = true;
+
+  return usable_devices;
+}
+
+struct GPUAddressHelper {
+  id<MTLBuffer> resource_buffer = nil;
+  id<MTLArgumentEncoder> address_encoder = nil;
+
+  /* One time setup of arg encoder. */
+  void init(id<MTLDevice> device)
+  {
+    if (resource_buffer) {
+      /* No setup required - already initialised. */
+      return;
+    }
+
+#  ifdef CYCLES_USE_TIER2D_BINDLESS
+    if (@available(macos 13.0, *)) {
+      /* No setup required - there's an API now! */
+      return;
+    }
+#  endif
+
+    /* Setup a tiny buffer to encode the GPU address / resourceID into. */
+    resource_buffer = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
+
+    /* Create an encoder to extract a gpuAddress from a MTLBuffer. */
+    MTLArgumentDescriptor *encoder_params = [[MTLArgumentDescriptor alloc] init];
+    encoder_params.arrayLength = 1;
+    encoder_params.access = MTLBindingAccessReadWrite;
+    encoder_params.dataType = MTLDataTypePointer;
+    address_encoder = [device newArgumentEncoderWithArguments:@[ encoder_params ]];
+    [address_encoder setArgumentBuffer:resource_buffer offset:0];
+  };
+
+  uint64_t gpuAddress(id<MTLBuffer> buffer)
+  {
+#  ifdef CYCLES_USE_TIER2D_BINDLESS
+    if (@available(macos 13.0, *)) {
+      return buffer.gpuAddress;
+    }
+#  endif
+    [address_encoder setBuffer:buffer offset:0 atIndex:0];
+    return *(uint64_t *)[resource_buffer contents];
+  }
+
+  uint64_t gpuResourceID(id<MTLTexture> texture)
+  {
+#  ifdef CYCLES_USE_TIER2D_BINDLESS
+    if (@available(macos 13.0, *)) {
+      MTLResourceID resourceID = texture.gpuResourceID;
+      return (uint64_t &)resourceID;
+    }
+#  endif
+    [address_encoder setTexture:texture atIndex:0];
+    return *(uint64_t *)[resource_buffer contents];
+  }
+
+  uint64_t gpuResourceID(id<MTLAccelerationStructure> accel_struct)
+  {
+#  ifdef CYCLES_USE_TIER2D_BINDLESS
+    if (@available(macos 13.0, *)) {
+      MTLResourceID resourceID = accel_struct.gpuResourceID;
+      return (uint64_t &)resourceID;
+    }
+#  endif
+    [address_encoder setAccelerationStructure:accel_struct atIndex:0];
+    return *(uint64_t *)[resource_buffer contents];
+  }
+
+  uint64_t gpuResourceID(id<MTLIntersectionFunctionTable> ift)
+  {
+#  ifdef CYCLES_USE_TIER2D_BINDLESS
+    if (@available(macos 13.0, *)) {
+      MTLResourceID resourceID = ift.gpuResourceID;
+      return (uint64_t &)resourceID;
+    }
+#  endif
+    [address_encoder setIntersectionFunctionTable:ift atIndex:0];
+    return *(uint64_t *)[resource_buffer contents];
+  }
+};
+
+GPUAddressHelper g_gpu_address_helper;
+
+void metal_gpu_address_helper_init(id<MTLDevice> device)
+{
+  g_gpu_address_helper.init(device);
+}
+
+uint64_t metal_gpuAddress(id<MTLBuffer> buffer)
+{
+  return g_gpu_address_helper.gpuAddress(buffer);
+}
+
+uint64_t metal_gpuResourceID(id<MTLTexture> texture)
+{
+  return g_gpu_address_helper.gpuResourceID(texture);
+}
+
+uint64_t metal_gpuResourceID(id<MTLAccelerationStructure> accel_struct)
+{
+  return g_gpu_address_helper.gpuResourceID(accel_struct);
+}
+
+uint64_t metal_gpuResourceID(id<MTLIntersectionFunctionTable> ift)
+{
+  return g_gpu_address_helper.gpuResourceID(ift);
+}
+
+CCL_NAMESPACE_END
+
+#endif /* WITH_METAL */