Asynchronously free ipiv 32-bit memory

src/lapack/backends/cusolver/cusolver_helper.hpp

@@ -284,6 +284,24 @@ inline int *create_dev_info(int num_elements = 1) {
     return reinterpret_cast<int *>(dev_info_d);
 }
 
+// Helper function for waiting on a vector of sycl events
+inline void depends_on_events(sycl::handler &cgh, std::vector<sycl::event> &dependencies = {}) {
+    for (sycl::event &e : dependencies)
+        cgh.depends_on(e);
+}
+
+// Asynchronously frees sycl USM `ptr` after waiting on events `dependencies`
+template <typename T>
+inline sycl::event free_async(sycl::queue &queue, T *ptr,
+                              std::vector<sycl::event> &dependencies = {}) {
+    sycl::event done = queue.submit([&](sycl::handler &cgh) {
+        depends_on_events(cgh, dependencies);
+
+        cgh.host_task([=](sycl::interop_handle ih) { sycl::free(ptr, queue); });
+    });
+    return done;
+}
+
 } // namespace cusolver
 } // namespace lapack
 } // namespace mkl

src/lapack/backends/cusolver/cusolver_lapack.cpp

@@ -1367,9 +1367,8 @@ inline sycl::event getrf(const char *func_name, Func func, sycl::queue &queue, s
         });
     });
 
-    queue.wait();
+    free_async(queue, ipiv32, { done_casting });
 
-    free(ipiv32, queue);
     return done_casting;
 }
 
@@ -1449,9 +1448,7 @@ inline sycl::event getrs(const char *func_name, Func func, sycl::queue &queue,
         });
     });
 
-    queue.wait();
-
-    free(ipiv32, queue);
+    free_async(queue, ipiv32, { done });
 
     return done;
 }
@@ -2207,9 +2204,8 @@ inline sycl::event sytrf(const char *func_name, Func func, sycl::queue &queue,
         });
     });
 
-    queue.wait();
+    free_async(queue, ipiv32, { done_casting });
 
-    free(ipiv32, queue);
     return done_casting;
 }