Fix ResidencySets reading autoreleased error object

mlx/backend/metal/resident.cpp

@@ -5,6 +5,7 @@
 #include <cstdio>
 #include <sstream>
 #include <stdexcept>
+#include <string>
 
 #include "mlx/backend/metal/device.h"
 #include "mlx/backend/metal/resident.h"
@@ -33,12 +34,12 @@ ResidencySets::ResidencySets(MTL::Device* d) {
     std::lock_guard<std::mutex> lk(mtx_);
     // Set 0 always exists and is the fallback when a later set cannot be
     // made, so failing to create it is fatal.
-    NS::Error* error = nullptr;
+    std::string error;
     if (!add_set_locked(&error)) {
       std::ostringstream msg;
       msg << "[metal::Device] Unable to construct residency set.\n";
-      if (error) {
-        msg << error->localizedDescription()->utf8String() << "\n";
+      if (!error.empty()) {
+        msg << error << "\n";
       }
       throw std::runtime_error(msg.str());
     }
@@ -47,7 +48,7 @@ ResidencySets::ResidencySets(MTL::Device* d) {
 
 ResidencySets::~ResidencySets() = default;
 
-bool ResidencySets::add_set_locked(NS::Error** error_out) {
+bool ResidencySets::add_set_locked(std::string* error_out) {
   NS::SharedPtr<MTL::ResidencySet> set;
   if (__builtin_available(macOS 15, iOS 18, *)) {
     auto pool = new_scoped_memory_pool();
@@ -58,8 +59,9 @@ bool ResidencySets::add_set_locked(NS::Error** error_out) {
       // A standing request, so allocations added to this set later are
       // covered without requesting residency again on every insert.
       set->requestResidency();
-    } else if (error_out) {
-      *error_out = error;
+    } else if (error_out && error) {
+      // The error is autoreleased into this pool, so copy its text out now.
+      *error_out = error->localizedDescription()->utf8String();
     }
   }
   if (!set) {

mlx/backend/metal/resident.h

@@ -5,6 +5,7 @@
 #include <atomic>
 #include <cstdint>
 #include <mutex>
+#include <string>
 #include <unordered_map>
 #include <vector>
 
@@ -96,7 +97,7 @@ class ResidencySets {
   // gave us a set; the new set is the last one. add_to_set_locked returns the
   // set it used. Neither it nor remove_from_set_locked commits, so a bulk
   // resize costs one commit per set instead of one per allocation.
-  bool add_set_locked(NS::Error** error = nullptr);
+  bool add_set_locked(std::string* error = nullptr);
   uint32_t choose_set_locked(size_t bytes);
   uint32_t add_to_set_locked(const MTL::Allocation* buf, Placement& at);
   void remove_from_set_locked(const MTL::Allocation* buf, Placement& at);