[SYCL][ABI-break] Remove the workaround for release of auxiliary buffers (#9914)

Remove the thread local variable workaround for deferred release of
auxiliary resources by passing the required information from reduction
headers.

Author: sergey-semenov

sycl/CMakeLists.txt

@@ -36,7 +36,7 @@ set(SYCL_MAJOR_VERSION 7)
 set(SYCL_MINOR_VERSION 0)
 set(SYCL_PATCH_VERSION 0)
 
-set(SYCL_DEV_ABI_VERSION 5)
+set(SYCL_DEV_ABI_VERSION 6)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)
   set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}")
 endif()

sycl/include/sycl/detail/cg.hpp

@@ -140,6 +140,12 @@ class CG {
   }
   std::vector<detail::EventImplPtr> &getEvents() { return MData.MEvents; }
 
+  virtual std::vector<std::shared_ptr<const void>>
+  getAuxiliaryResources() const {
+    return {};
+  }
+  virtual void clearAuxiliaryResources(){};
+
   virtual ~CG() = default;
 
 private:
@@ -198,32 +204,40 @@ class CGExecKernel : public CG {
     return MStreams;
   }
 
-  std::vector<std::shared_ptr<const void>> getAuxiliaryResources() const {
+  std::vector<std::shared_ptr<const void>>
+  getAuxiliaryResources() const override {
     return MAuxiliaryResources;
   }
+  void clearAuxiliaryResources() override { MAuxiliaryResources.clear(); }
 
   std::shared_ptr<detail::kernel_bundle_impl> getKernelBundle() {
     return MKernelBundle;
   }
 
   void clearStreams() { MStreams.clear(); }
   bool hasStreams() { return !MStreams.empty(); }
-
-  void clearAuxiliaryResources() { MAuxiliaryResources.clear(); }
-  bool hasAuxiliaryResources() { return !MAuxiliaryResources.empty(); }
 };
 
 /// "Copy memory" command group class.
 class CGCopy : public CG {
   void *MSrc;
   void *MDst;
+  std::vector<std::shared_ptr<const void>> MAuxiliaryResources;
 
 public:
   CGCopy(CGTYPE CopyType, void *Src, void *Dst, CG::StorageInitHelper CGData,
+         std::vector<std::shared_ptr<const void>> AuxiliaryResources,
          detail::code_location loc = {})
-      : CG(CopyType, std::move(CGData), std::move(loc)), MSrc(Src), MDst(Dst) {}
+      : CG(CopyType, std::move(CGData), std::move(loc)), MSrc(Src), MDst(Dst),
+        MAuxiliaryResources{AuxiliaryResources} {}
   void *getSrc() { return MSrc; }
   void *getDst() { return MDst; }
+
+  std::vector<std::shared_ptr<const void>>
+  getAuxiliaryResources() const override {
+    return MAuxiliaryResources;
+  }
+  void clearAuxiliaryResources() override { MAuxiliaryResources.clear(); }
 };
 
 /// "Fill memory" command group class.

sycl/include/sycl/detail/helpers.hpp

@@ -37,6 +37,7 @@ enum class memory_order;
 
 namespace detail {
 
+class buffer_impl;
 class context_impl;
 // The function returns list of events that can be passed to OpenCL API as
 // dependency list and waits for others.
@@ -46,6 +47,9 @@ getOrWaitEvents(std::vector<sycl::event> DepEvents,
 
 __SYCL_EXPORT void waitEvents(std::vector<sycl::event> DepEvents);
 
+__SYCL_EXPORT void
+markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl);
+
 template <typename T> T *declptr() { return static_cast<T *>(nullptr); }
 
 // Function to get of store id, item, nd_item, group for the host implementation
@@ -246,7 +250,6 @@ void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
 template <size_t count, class F> void loop(F &&f) {
   loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
 }
-
 } // namespace detail
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

sycl/include/sycl/handler.hpp

@@ -457,13 +457,26 @@ class __SYCL_EXPORT handler {
     MStreamStorage.push_back(Stream);
   }
 
-  /// Saves buffers created by handling reduction feature in handler.
+  /// Saves resources created by handling reduction feature in handler.
   /// They are then forwarded to command group and destroyed only after
   /// the command group finishes the work on device/host.
   ///
   /// @param ReduObj is a pointer to object that must be stored.
   void addReduction(const std::shared_ptr<const void> &ReduObj);
 
+  /// Saves buffers created by handling reduction feature in handler and marks
+  /// them as internal. They are then forwarded to command group and destroyed
+  /// only after the command group finishes the work on device/host.
+  ///
+  /// @param ReduBuf is a pointer to buffer that must be stored.
+  template <typename T, int Dimensions, typename AllocatorT, typename Enable>
+  void
+  addReduction(const std::shared_ptr<buffer<T, Dimensions, AllocatorT, Enable>>
+                   &ReduBuf) {
+    detail::markBufferAsInternal(getSyclObjImpl(*ReduBuf));
+    addReduction(std::shared_ptr<const void>(ReduBuf));
+  }
+
   ~handler() = default;
 
   // TODO: Private and unusued. Remove when ABI break is allowed.

sycl/include/sycl/reduction.hpp

@@ -960,11 +960,9 @@ class reduction_impl_algo {
       for (int i = 0; i < num_elements; ++i) {
         (*RWReduVal)[i] = decltype(MIdentityContainer)::getIdentity();
       }
-      CGH.addReduction(RWReduVal);
       auto Buf = std::make_shared<buffer<T, 1>>(RWReduVal.get()->data(),
                                                 range<1>(num_elements));
       Buf->set_final_data();
-      CGH.addReduction(Buf);
       accessor Mem{*Buf, CGH};
       Func(Mem);
 
@@ -975,6 +973,10 @@ class reduction_impl_algo {
         // so use the old-style API.
         auto Mem =
             Buf->template get_access<access::mode::read_write>(CopyHandler);
+        // Since this CG is dependent on the one associated with CGH,
+        // registering the auxiliary resources here is enough.
+        CopyHandler.addReduction(RWReduVal);
+        CopyHandler.addReduction(Buf);
         if constexpr (is_usm) {
           // Can't capture whole reduction, copy into distinct variables.
           bool IsUpdateOfUserVar = !initializeToIdentity();

sycl/source/detail/helpers.cpp

@@ -9,6 +9,7 @@
 #include <detail/scheduler/commands.hpp>
 #include <sycl/detail/helpers.hpp>
 
+#include <detail/buffer_impl.hpp>
 #include <detail/context_impl.hpp>
 #include <detail/event_impl.hpp>
 #include <detail/queue_impl.hpp>
@@ -63,6 +64,10 @@ void waitEvents(std::vector<sycl::event> DepEvents) {
   }
 }
 
+void markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl) {
+  BufImpl->markAsInternal();
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/detail/scheduler/scheduler.cpp

@@ -93,14 +93,11 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
   const CG::CGTYPE Type = CommandGroup->getType();
   std::vector<Command *> AuxiliaryCmds;
   std::vector<StreamImplPtr> Streams;
-  std::vector<std::shared_ptr<const void>> AuxiliaryResources;
 
   if (Type == CG::Kernel) {
     auto *CGExecKernelPtr = static_cast<CGExecKernel *>(CommandGroup.get());
     Streams = CGExecKernelPtr->getStreams();
     CGExecKernelPtr->clearStreams();
-    AuxiliaryResources = CGExecKernelPtr->getAuxiliaryResources();
-    CGExecKernelPtr->clearAuxiliaryResources();
     // Stream's flush buffer memory is mainly initialized in stream's __init
     // method. However, this method is not available on host device.
     // Initializing stream's flush buffer on the host side in a separate task.
@@ -110,6 +107,9 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
       }
     }
   }
+  std::vector<std::shared_ptr<const void>> AuxiliaryResources;
+  AuxiliaryResources = CommandGroup->getAuxiliaryResources();
+  CommandGroup->clearAuxiliaryResources();
 
   bool ShouldEnqueue = true;
   {
@@ -546,7 +546,6 @@ void Scheduler::registerAuxiliaryResources(
 
 void Scheduler::cleanupAuxiliaryResources(BlockingT Blocking) {
   std::unique_lock<std::mutex> Lock{MAuxiliaryResourcesMutex};
-  ForceDeferredReleaseWrapper ForceDeferredRelease;
   for (auto It = MAuxiliaryResources.begin();
        It != MAuxiliaryResources.end();) {
     const EventImplPtr &Event = It->first;
@@ -560,8 +559,6 @@ void Scheduler::cleanupAuxiliaryResources(BlockingT Blocking) {
   }
 }
 
-thread_local bool Scheduler::ForceDeferredMemObjRelease = false;
-
 void Scheduler::startFusion(QueueImplPtr Queue) {
   WriteLockT Lock = acquireWriteLock();
   MGraphBuilder.startFusion(Queue);

sycl/source/detail/scheduler/scheduler.hpp

@@ -870,21 +870,6 @@ class Scheduler {
 
   QueueImplPtr DefaultHostQueue;
 
-  // This thread local flag is a workaround for a problem with managing
-  // auxiliary resources. We would like to release internal buffers used for
-  // reductions in a deferred manner, but marking them individually isn't an
-  // option since all auxiliary resources (buffers, host memory, USM) are passed
-  // to the library as type erased shared pointers. This flag makes it so that
-  // release of every memory object is deferred while it's set, and it should
-  // only be set during release of auxiliary resources.
-  // TODO Remove once ABI breaking changes are allowed.
-  friend class SYCLMemObjT;
-  static thread_local bool ForceDeferredMemObjRelease;
-  struct ForceDeferredReleaseWrapper {
-    ForceDeferredReleaseWrapper() { ForceDeferredMemObjRelease = true; };
-    ~ForceDeferredReleaseWrapper() { ForceDeferredMemObjRelease = false; };
-  };
-
   friend class Command;
   friend class DispatchHostTask;
   friend class queue_impl;

sycl/source/detail/sycl_mem_obj_t.cpp

@@ -219,10 +219,7 @@ void SYCLMemObjT::detachMemoryObject(
   // buffer creation and set to meaningfull
   // value only if any operation on buffer submitted inside addCG call. addCG is
   // called from queue::submit and buffer destruction could not overlap with it.
-  // ForceDeferredMemObjRelease is a workaround for managing auxiliary resources
-  // while preserving backward compatibility, see the comment for
-  // ForceDeferredMemObjRelease in scheduler.
-  if (MRecord && (!MHostPtrProvided || Scheduler::ForceDeferredMemObjRelease))
+  if (MRecord && (!MHostPtrProvided || MIsInternal))
     Scheduler::getInstance().deferMemObjRelease(Self);
 }
 

sycl/source/detail/sycl_mem_obj_t.hpp

@@ -272,6 +272,8 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
 
   void detachMemoryObject(const std::shared_ptr<SYCLMemObjT> &Self) const;
 
+  void markAsInternal() { MIsInternal = true; }
+
 protected:
   // An allocateMem helper that determines which host ptr to use
   void determineHostPtr(const ContextImplPtr &Context, bool InitFromUserData,
@@ -312,6 +314,10 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
   // we have read only HostPtr - MUploadDataFunctor is empty but delayed release
   // must be not allowed.
   bool MHostPtrProvided;
+  // Indicates that the memory object was allocated internally. Such memory
+  // objects can be released in a deferred manner regardless of whether a host
+  // pointer was provided or not.
+  bool MIsInternal = false;
 };
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

sycl/source/handler.cpp

@@ -293,8 +293,9 @@ event handler::finalize() {
   case detail::CG::CopyAccToPtr:
   case detail::CG::CopyPtrToAcc:
   case detail::CG::CopyAccToAcc:
-    CommandGroup.reset(new detail::CGCopy(MCGType, MSrcPtr, MDstPtr,
-                                          std::move(CGData), MCodeLoc));
+    CommandGroup.reset(
+        new detail::CGCopy(MCGType, MSrcPtr, MDstPtr, std::move(CGData),
+                           std::move(MImpl->MAuxiliaryResources), MCodeLoc));
     break;
   case detail::CG::Fill:
     CommandGroup.reset(new detail::CGFill(std::move(MPattern), MDstPtr,

sycl/test/abi/sycl_symbols_linux.dump

@@ -3896,6 +3896,7 @@ _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostE
 _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_28SampledImageAccessorBaseHostENS0_12image_targetE
 _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_30UnsampledImageAccessorBaseHostENS0_12image_targetE
 _ZN4sycl3_V16detail20getDeviceFromHandlerERNS0_7handlerE
+_ZN4sycl3_V16detail20markBufferAsInternalERKSt10shared_ptrINS1_11buffer_implEE
 _ZN4sycl3_V16detail21LocalAccessorBaseHost12getNumOfDimsEv
 _ZN4sycl3_V16detail21LocalAccessorBaseHost14getElementSizeEv
 _ZN4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv

sycl/test/abi/sycl_symbols_windows.dump

@@ -1242,6 +1242,8 @@
 ?malloc_shared@_V1@sycl@@YAPEAX_KAEBVqueue@12@AEBUcode_location@detail@12@@Z
 ?malloc_shared@_V1@sycl@@YAPEAX_KAEBVqueue@12@AEBVproperty_list@12@AEBUcode_location@detail@12@@Z
 ?map@MemoryManager@detail@_V1@sycl@@SAPEAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@W4mode@access@34@IV?$range@$02@34@4V?$id@$02@34@IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z
+?markAsInternal@SYCLMemObjT@detail@_V1@sycl@@QEAAXXZ
+?markBufferAsInternal@detail@_V1@sycl@@YAXAEBV?$shared_ptr@Vbuffer_impl@detail@_V1@sycl@@@std@@@Z
 ?mem_advise@handler@_V1@sycl@@QEAAXPEBX_KH@Z
 ?mem_advise@queue@_V1@sycl@@QEAA?AVevent@23@PEBX_KHAEBUcode_location@detail@23@@Z
 ?mem_advise@queue@_V1@sycl@@QEAA?AVevent@23@PEBX_KHAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z

sycl/test/abi/vtable.cpp

@@ -24,12 +24,13 @@ void foo(sycl::detail::HostKernelBase &HKB) {
 // CHECK-NEXT:   5 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting]
 
 void foo(sycl::detail::CG *CG) { delete CG; }
-// CHECK:    Vtable for 'sycl::detail::CG' (4 entries).
-// CHECK-NEXT:   0 | offset_to_top (0)
-// CHECK-NEXT:   1 | sycl::detail::CG RTTI
-// CHECK-NEXT:       -- (sycl::detail::CG, 0) vtable address --
-// CHECK-NEXT:   2 | sycl::detail::CG::~CG() [complete]
-// CHECK-NEXT:   3 | sycl::detail::CG::~CG() [deleting]
+// CHECK: Vtable for 'sycl::detail::CG' (6 entries).
+// CHECK-NEXT:    0 | offset_to_top (0)
+// CHECK-NEXT:    1 | sycl::detail::CG RTTI
+// CHECK-NEXT:        -- (sycl::detail::CG, 0) vtable address --
+// CHECK-NEXT:    2 | std::vector<std::shared_ptr<const void>> sycl::detail::CG::getAuxiliaryResources() const
+// CHECK-NEXT:    3 | void sycl::detail::CG::clearAuxiliaryResources()
+// CHECK-NEXT:    4 | sycl::detail::CG::~CG() [complete]
 
 void foo(sycl::detail::PropertyWithDataBase *Prop) { delete Prop; }
 // CHECK:    Vtable for 'sycl::detail::PropertyWithDataBase' (4 entries).