[SYCL][COMPAT] Add wait_and_free plus rename async_free in syclcompat (…

…intel#14015)

This PR adds a `wait_and_free` func. This makes it safer and less likely
to release memory during or before it is used by enqueued commands.

`async_free` is renamed `enqueue_free`, to make its behaviour clearer

This PR updates the comments and tests accordingly

sycl/doc/syclcompat/README.md

@@ -489,10 +489,12 @@ sycl::event memset_async(pitched_data pitch, int val,
                          sycl::range<3> size,
                          sycl::queue q = get_default_queue()); // 3D matrix
 
+// Free
+void wait_and_free(void *ptr, sycl::queue q = get_default_queue());
 void free(void *ptr, sycl::queue q = get_default_queue());
-sycl::event free_async(const std::vector<void *> &pointers,
-                       const std::vector<sycl::event> &events,
-                       sycl::queue q = get_default_queue());
+sycl::event enqueue_free(const std::vector<void *> &pointers,
+                         const std::vector<sycl::event> &events,
+                         sycl::queue q = get_default_queue());
 
 // Queries pointer allocation type
 class pointer_attributes {
@@ -870,7 +872,7 @@ independently of what is set in this parameter.
 Devices are managed through a helper class, `device_ext`. The `device_ext` class
 associates a vector of `sycl::queues` with its `sycl::device`. The `device_ext`
 destructor waits on a set of `sycl::event` which can be added to via
-`add_event`. This is used, for example, to implement `syclcompat::free_async` to
+`add_event`. This is used, for example, to implement `syclcompat::enqueue_free` to
 schedule release of memory after a kernel or `mempcy`. SYCL device properties
 can be queried through `device_ext` as well.
 `device_ext` also provides the `has_capability_or_fail` member function, which

sycl/include/syclcompat/device.hpp

@@ -664,8 +664,9 @@ Use 64 bits as memory_bus_width default value."
     std::lock_guard<std::mutex> lock(m_mutex);
     _events.push_back(event);
   }
-  friend sycl::event free_async(const std::vector<void *> &,
-                                const std::vector<sycl::event> &, sycl::queue);
+  friend sycl::event enqueue_free(const std::vector<void *> &,
+                                  const std::vector<sycl::event> &,
+                                  sycl::queue);
   queue_ptr _default_queue;
   queue_ptr _saved_queue;
   sycl::context _ctx;

sycl/include/syclcompat/memory.hpp

@@ -531,26 +531,40 @@ static inline void *malloc(size_t &pitch, size_t x, size_t y,
   return detail::malloc(pitch, x, y, 1, q);
 }
 
-/// free
+/// Wait on the queue \p q and free the memory \p ptr.
 /// \param ptr Point to free.
 /// \param q Queue to execute the free task.
 /// \returns no return value.
+static inline void wait_and_free(void *ptr,
+                                 sycl::queue q = get_default_queue()) {
+  get_current_device().queues_wait_and_throw();
+  q.wait();
+  if (ptr) {
+    sycl::free(ptr, q);
+  }
+}
+
+/// Free the memory \p ptr on the default queue without synchronizing
+/// \param ptr Point to free.
+/// \returns no return value.
 static inline void free(void *ptr, sycl::queue q = get_default_queue()) {
   if (ptr) {
-    sycl::free(ptr, q.get_context());
+    sycl::free(ptr, q);
   }
 }
 
-/// Free the device memory pointed by a batch of pointers in \p pointers which
-/// are related to \p q after \p events completed.
+/// Enqueues the release of all pointers in /p pointers on the /p q.
+/// The command waits on all passed /p events and returns an event that
+/// track the commands execution on the queue.
 ///
 /// \param pointers The pointers point to the device memory requested to be
-/// freed. \param events The events to be waited. \param q The sycl::queue the
-/// memory relates to.
+/// freed.
+/// \param events The events to be waited on.
+/// \param q The sycl::queue the memory relates to.
 // Can't be static due to the friend declaration in the memory header.
-inline sycl::event free_async(const std::vector<void *> &pointers,
-                              const std::vector<sycl::event> &events,
-                              sycl::queue q = get_default_queue()) {
+inline sycl::event enqueue_free(const std::vector<void *> &pointers,
+                                const std::vector<sycl::event> &events,
+                                sycl::queue q = get_default_queue()) {
   auto event = q.submit(
       [&pointers, &events, ctxt = q.get_context()](sycl::handler &cgh) {
         cgh.depends_on(events);

sycl/test-e2e/syclcompat/memory/memory_async.cpp

@@ -43,14 +43,15 @@
 
 #include "memory_fixt.hpp"
 
-// free_async is a host task, so we are really testing the event dependency here
+// enqueue_free is just a host task, so we are really testing the event
+// dependency here
 void test_free_async() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
   AsyncTest atest;
 
   float *d_D = (float *)syclcompat::malloc(sizeof(float));
   sycl::event kernel_ev = atest.launch_kernel();
-  sycl::event free_ev = syclcompat::free_async({d_D}, {kernel_ev});
+  sycl::event free_ev = syclcompat::enqueue_free({d_D}, {kernel_ev});
 
   atest.check_events(kernel_ev, free_ev);
 }

sycl/test-e2e/syclcompat/memory/memory_management_test3.cpp

@@ -65,6 +65,30 @@ void test_free_memory_q() {
   syclcompat::free(nullptr, q);
 }
 
+void test_wait_and_free_memory() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  sycl::queue q{{sycl::property::queue::in_order()}};
+  float *d_A = (float *)syclcompat::malloc(sizeof(float), q);
+  syclcompat::wait_and_free((void *)d_A);
+
+  syclcompat::wait_and_free(0);
+  syclcompat::wait_and_free(NULL);
+  syclcompat::wait_and_free(nullptr);
+}
+
+void test_wait_and_free_memory_q() {
+  std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+  sycl::queue q{{sycl::property::queue::in_order()}};
+  float *d_A = (float *)syclcompat::malloc(sizeof(float), q);
+  syclcompat::wait_and_free((void *)d_A, q);
+
+  syclcompat::wait_and_free(0, q);
+  syclcompat::wait_and_free(NULL, q);
+  syclcompat::wait_and_free(nullptr, q);
+}
+
 void test_memcpy_async() {
   std::cout << __PRETTY_FUNCTION__ << std::endl;
 
@@ -662,6 +686,8 @@ void test_constant_memcpy_async_q() {
 int main() {
   test_free_memory();
   test_free_memory_q();
+  test_wait_and_free_memory();
+  test_wait_and_free_memory_q();
   test_memcpy_async();
   test_memcpy_async_q();
   test_memcpy_async_pitched();