[SYCL] Fix resource leak related to SYCL_FALLBACK_ASSERT (#12532)

#6837 enabled asynchronous buffer
destruction for buffers constructed without host data. However, initial
fallback assert implementation in
#3767 predates it and as such had to
place the buffer inside `queue_impl` to avoid unintended synchronization
point. I don't know if there was the same crash observed on the
end-to-end test added as part of this PR prior to
#3767, but it doesn't even matter
because the "new" implementation is both simpler and doesn't result in a
crash.

I suspect that without it (with the buffer for fallback assert
implementation being a data member of `sycl::queue_impl`) we had a
cyclic dependency somewhere leading to resource leak and ultimately to
the assert in `DeviceGlobalUSMMem::~DeviceGlobalUSMMem()`.

Author: aelovikov-intel

sycl/include/sycl/queue.hpp

@@ -2965,7 +2965,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
         Rest...);
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   buffer<detail::AssertHappened, 1> &getAssertHappenedBuffer();
+#endif
 
   event memcpyToDeviceGlobal(void *DeviceGlobalPtr, const void *Src,
                              bool IsDeviceImageScope, size_t NumBytes,
@@ -3019,9 +3021,7 @@ class AssertInfoCopier;
  */
 event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
                           const detail::code_location &CodeLoc) {
-  using AHBufT = buffer<detail::AssertHappened, 1>;
-
-  AHBufT &Buffer = Self.getAssertHappenedBuffer();
+  buffer<detail::AssertHappened, 1> Buffer{1};
 
   event CopierEv, CheckerEv, PostCheckerEv;
   auto CopierCGF = [&](handler &CGH) {

sycl/source/detail/queue_impl.hpp

@@ -108,7 +108,9 @@ class queue_impl {
              const async_handler &AsyncHandler, const property_list &PropList)
       : MDevice(Device), MContext(Context), MAsyncHandler(AsyncHandler),
         MPropList(PropList), MHostQueue(MDevice->is_host()),
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
         MAssertHappenedBuffer(range<1>{1}),
+#endif
         MIsInorder(has_property<property::queue::in_order>()),
         MDiscardEvents(
             has_property<ext::oneapi::property::queue::discard_events>()),
@@ -283,7 +285,9 @@ class queue_impl {
   queue_impl(sycl::detail::pi::PiQueue PiQueue, const ContextImplPtr &Context,
              const async_handler &AsyncHandler)
       : MContext(Context), MAsyncHandler(AsyncHandler), MHostQueue(false),
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
         MAssertHappenedBuffer(range<1>{1}),
+#endif
         MIsInorder(has_property<property::queue::in_order>()),
         MDiscardEvents(
             has_property<ext::oneapi::property::queue::discard_events>()),
@@ -305,7 +309,10 @@ class queue_impl {
   queue_impl(sycl::detail::pi::PiQueue PiQueue, const ContextImplPtr &Context,
              const async_handler &AsyncHandler, const property_list &PropList)
       : MContext(Context), MAsyncHandler(AsyncHandler), MPropList(PropList),
-        MHostQueue(false), MAssertHappenedBuffer(range<1>{1}),
+        MHostQueue(false),
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+        MAssertHappenedBuffer(range<1>{1}),
+#endif
         MIsInorder(has_property<property::queue::in_order>()),
         MDiscardEvents(
             has_property<ext::oneapi::property::queue::discard_events>()),
@@ -673,9 +680,11 @@ class queue_impl {
   /// \return a native handle.
   pi_native_handle getNative(int32_t &NativeHandleDesc) const;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   buffer<AssertHappened, 1> &getAssertHappenedBuffer() {
     return MAssertHappenedBuffer;
   }
+#endif
 
   void registerStreamServiceEvent(const EventImplPtr &Event) {
     std::lock_guard<std::mutex> Lock(MMutex);
@@ -918,8 +927,10 @@ class queue_impl {
   /// need to emulate it with multiple native in-order queues.
   bool MEmulateOOO = false;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // Buffer to store assert failure descriptor
   buffer<AssertHappened, 1> MAssertHappenedBuffer;
+#endif
 
   // This event is employed for enhanced dependency tracking with in-order queue
   // Access to the event should be guarded with MMutex

sycl/source/queue.cpp

@@ -298,9 +298,11 @@ pi_native_handle queue::getNative(int32_t &NativeHandleDesc) const {
   return impl->getNative(NativeHandleDesc);
 }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 buffer<detail::AssertHappened, 1> &queue::getAssertHappenedBuffer() {
   return impl->getAssertHappenedBuffer();
 }
+#endif
 
 event queue::memcpyToDeviceGlobal(void *DeviceGlobalPtr, const void *Src,
                                   bool IsDeviceImageScope, size_t NumBytes,

sycl/test-e2e/Assert/check_resource_leak.cpp

@@ -0,0 +1,32 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Device globals aren't supported on opencl:gpu yet.
+// UNSUPPORTED: opencl && gpu
+
+// TODO: Fails at JIT compilation for some reason.
+// UNSUPPORTED: hip
+#define SYCL_FALLBACK_ASSERT 1
+
+#include <sycl/sycl.hpp>
+
+// DeviceGlobalUSMMem::~DeviceGlobalUSMMem() has asserts to ensure some
+// resources have been cleaned up when it's executed. Those asserts used to fail
+// when "AssertHappened" buffer used in fallback implementation of the device
+// assert was a data member of the queue_impl.
+sycl::ext::oneapi::experimental::device_global<int32_t> dg;
+
+int main() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &cgh) {
+     sycl::range<1> R{16};
+     cgh.parallel_for(sycl::nd_range<1>{R, R}, [=](sycl::nd_item<1> ndi) {
+       if (ndi.get_global_linear_id() == 0)
+         dg.get() = 42;
+       auto sg = sycl::ext::oneapi::experimental::this_sub_group();
+       auto active = sycl::ext::oneapi::group_ballot(sg, 1);
+     });
+   }).wait();
+
+  return 0;
+}