[SYCL][NFC] Have separate code path for secondary queue + remove duplicate parameters (#18201)

Signed-off-by: Agarwal, Udit <[email protected]>

Author: uditagarwal97

sycl/include/sycl/handler.hpp

@@ -448,7 +448,6 @@ class __SYCL_EXPORT handler {
           bool CallerNeedsEvent);
 #endif
   __SYCL_DLL_LOCAL handler(std::shared_ptr<detail::queue_impl> Queue,
-                           detail::queue_impl *PrimaryQueue,
                            detail::queue_impl *SecondaryQueue,
                            bool CallerNeedsEvent);
 

sycl/include/sycl/queue.hpp

@@ -373,8 +373,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
       const detail::code_location &CodeLoc = detail::code_location::current()) {
     return submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
         sycl::ext::oneapi::experimental::empty_properties_t{},
-        detail::type_erased_cgfo_ty{CGF},
-        /*SecondaryQueuePtr=*/nullptr, CodeLoc);
+        detail::type_erased_cgfo_ty{CGF}, CodeLoc);
   }
 
   /// Submits a command group function object to the queue, in order to be
@@ -3609,6 +3608,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   ///
   /// \param Props is a property list with submission properties.
   /// \param CGF is a function object containing command group.
+  /// \param SecondaryQueuePtr is a pointer to the secondary queue.
   /// \param CodeLoc is the code location of the submit call (default argument)
   /// \return a SYCL event object for the submitted command group.
   //
@@ -3643,6 +3643,41 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                   TlsCodeLocCapture.isToplevel());
   }
 
+  /// Submits a command group function object to the queue, in order to be
+  /// scheduled for execution on the device.
+  ///
+  /// \param Props is a property list with submission properties.
+  /// \param CGF is a function object containing command group.
+  /// \param CodeLoc is the code location of the submit call (default argument)
+  /// \return a SYCL event object for the submitted command group.
+  //
+  // UseFallBackAssert as template param vs `#if` in function body is necessary
+  // to prevent ODR-violation between TUs built with different fallback assert
+  // modes.
+  template <bool UseFallbackAssert, typename PropertiesT>
+  event submit_with_event(
+      PropertiesT Props, const detail::type_erased_cgfo_ty &CGF,
+      const detail::code_location &CodeLoc = detail::code_location::current()) {
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+    detail::SubmissionInfo SI{};
+    ProcessSubmitProperties(Props, SI);
+    if constexpr (UseFallbackAssert)
+      SI.PostProcessorFunc() = [this, &TlsCodeLocCapture](bool IsKernel,
+                                                          bool KernelUsesAssert,
+                                                          event &E) {
+        if (IsKernel && !device_has(aspect::ext_oneapi_native_assert) &&
+            KernelUsesAssert && !device_has(aspect::accelerator)) {
+          // __devicelib_assert_fail isn't supported by Device-side Runtime
+          // Linking against fallback impl of __devicelib_assert_fail is
+          // performed by program manager class
+          // Fallback assert isn't supported for FPGA
+          submitAssertCapture(*this, E, nullptr, TlsCodeLocCapture.query());
+        }
+      };
+    return submit_with_event_impl(CGF, SI, TlsCodeLocCapture.query(),
+                                  TlsCodeLocCapture.isToplevel());
+  }
+
   /// Submits a command group function object to the queue, in order to be
   /// scheduled for execution on the device.
   ///
@@ -3660,7 +3695,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     if constexpr (UseFallbackAssert) {
       // If post-processing is needed, fall back to the regular submit.
       // TODO: Revisit whether we can avoid this.
-      submit_with_event<UseFallbackAssert>(Props, CGF, nullptr, CodeLoc);
+      submit_with_event<UseFallbackAssert>(Props, CGF, CodeLoc);
     } else {
       detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
       detail::SubmissionInfo SI{};

sycl/source/detail/handler_impl.hpp

@@ -31,10 +31,8 @@ enum class HandlerSubmissionState : std::uint8_t {
 
 class handler_impl {
 public:
-  handler_impl(queue_impl *SubmissionPrimaryQueue,
-               queue_impl *SubmissionSecondaryQueue, bool EventNeeded)
-      : MSubmissionPrimaryQueue(SubmissionPrimaryQueue),
-        MSubmissionSecondaryQueue(SubmissionSecondaryQueue),
+  handler_impl(queue_impl *SubmissionSecondaryQueue, bool EventNeeded)
+      : MSubmissionSecondaryQueue(SubmissionSecondaryQueue),
         MEventNeeded(EventNeeded) {};
 
   handler_impl(
@@ -69,15 +67,8 @@ class handler_impl {
   /// Registers mutually exclusive submission states.
   HandlerSubmissionState MSubmissionState = HandlerSubmissionState::NO_STATE;
 
-  /// Shared pointer to the primary queue implementation. This is different from
-  /// the queue associated with the handler if the corresponding submission is
-  /// a fallback from a previous submission.
-  queue_impl *MSubmissionPrimaryQueue = nullptr;
-
-  /// Shared pointer to the secondary queue implementation. Nullptr if no
-  /// secondary queue fallback was given in the associated submission. This is
-  /// equal to the queue associated with the handler if the corresponding
-  /// submission is a fallback from a previous submission.
+  /// Pointer to the secondary queue implementation. Nullptr if no
+  /// secondary queue fallback was given in the associated submission.
   queue_impl *MSubmissionSecondaryQueue = nullptr;
 
   /// Bool stores information about whether the event resulting from the

sycl/source/detail/queue_impl.cpp

@@ -310,6 +310,57 @@ void queue_impl::addEvent(const event &Event) {
   }
 }
 
+event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
+                              const std::shared_ptr<queue_impl> &Self,
+                              queue_impl *SecondaryQueue, bool CallerNeedsEvent,
+                              const detail::code_location &Loc,
+                              bool IsTopCodeLoc,
+                              const SubmissionInfo &SubmitInfo) {
+  handler Handler(Self, SecondaryQueue, CallerNeedsEvent);
+  auto &HandlerImpl = detail::getSyclObjImpl(Handler);
+  if (xptiTraceEnabled()) {
+    Handler.saveCodeLoc(Loc, IsTopCodeLoc);
+  }
+
+  {
+    NestedCallsTracker tracker;
+    CGF(Handler);
+  }
+
+  // Scheduler will later omit events, that are not required to execute tasks.
+  // Host and interop tasks, however, are not submitted to low-level runtimes
+  // and require separate dependency management.
+  const CGType Type = HandlerImpl->MCGType;
+  std::vector<StreamImplPtr> Streams;
+  if (Type == CGType::Kernel)
+    Streams = std::move(Handler.MStreamStorage);
+
+  HandlerImpl->MEventMode = SubmitInfo.EventMode();
+
+  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc());
+
+  addEvent(Event);
+
+  const auto &EventImpl = detail::getSyclObjImpl(Event);
+  for (auto &Stream : Streams) {
+    // We don't want stream flushing to be blocking operation that is why submit
+    // a host task to print stream buffer. It will fire up as soon as the kernel
+    // finishes execution.
+    auto L = [&](handler &ServiceCGH) {
+      Stream->generateFlushCommand(ServiceCGH);
+    };
+    detail::type_erased_cgfo_ty CGF{L};
+    event FlushEvent =
+        submit_impl(CGF, Self, SecondaryQueue, /*CallerNeedsEvent*/ true, Loc,
+                    IsTopCodeLoc, {});
+    EventImpl->attachEventToCompleteWeak(detail::getSyclObjImpl(FlushEvent));
+    registerStreamServiceEvent(detail::getSyclObjImpl(FlushEvent));
+  }
+
+  return Event;
+}
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
                               const std::shared_ptr<queue_impl> &Self,
                               const std::shared_ptr<queue_impl> &PrimaryQueue,
@@ -318,8 +369,7 @@ event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
                               const detail::code_location &Loc,
                               bool IsTopCodeLoc,
                               const SubmissionInfo &SubmitInfo) {
-  handler Handler(Self, PrimaryQueue.get(), SecondaryQueue.get(),
-                  CallerNeedsEvent);
+  handler Handler(Self, CallerNeedsEvent);
   auto &HandlerImpl = detail::getSyclObjImpl(Handler);
 
 #if XPTI_ENABLE_INSTRUMENTATION
@@ -365,6 +415,7 @@ event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
 
   return Event;
 }
+#endif
 
 template <typename HandlerFuncT>
 event queue_impl::submitWithHandler(const std::shared_ptr<queue_impl> &Self,

sycl/source/detail/queue_impl.hpp

@@ -376,7 +376,7 @@ class queue_impl {
                           const detail::code_location &Loc, bool IsTopCodeLoc) {
 
     event ResEvent =
-        submit_impl(CGF, Self, Self, SubmitInfo.SecondaryQueue(),
+        submit_impl(CGF, Self, SubmitInfo.SecondaryQueue().get(),
                     /*CallerNeedsEvent=*/true, Loc, IsTopCodeLoc, SubmitInfo);
     return discard_or_return(ResEvent);
   }
@@ -386,7 +386,7 @@ class queue_impl {
                             const SubmissionInfo &SubmitInfo,
                             const detail::code_location &Loc,
                             bool IsTopCodeLoc) {
-    submit_impl(CGF, Self, Self, SubmitInfo.SecondaryQueue(),
+    submit_impl(CGF, Self, SubmitInfo.SecondaryQueue().get(),
                 /*CallerNeedsEvent=*/false, Loc, IsTopCodeLoc, SubmitInfo);
   }
 
@@ -840,6 +840,7 @@ class queue_impl {
     }
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   /// Performs command group submission to the queue.
   ///
   /// \param CGF is a function object containing command group.
@@ -859,6 +860,23 @@ class queue_impl {
                     const std::shared_ptr<queue_impl> &SecondaryQueue,
                     bool CallerNeedsEvent, const detail::code_location &Loc,
                     bool IsTopCodeLoc, const SubmissionInfo &SubmitInfo);
+#endif
+
+  /// Performs command group submission to the queue.
+  ///
+  /// \param CGF is a function object containing command group.
+  /// \param Self is a pointer to this queue.
+  /// \param SecondaryQueue is a pointer to the secondary queue.
+  /// \param CallerNeedsEvent is a boolean indicating whether the event is
+  ///        required by the user after the call.
+  /// \param Loc is the code location of the submit call (default argument)
+  /// \param SubmitInfo is additional optional information for the submission.
+  /// \return a SYCL event representing submitted command group.
+  event submit_impl(const detail::type_erased_cgfo_ty &CGF,
+                    const std::shared_ptr<queue_impl> &Self,
+                    queue_impl *SecondaryQueue, bool CallerNeedsEvent,
+                    const detail::code_location &Loc, bool IsTopCodeLoc,
+                    const SubmissionInfo &SubmitInfo);
 
   /// Helper function for submitting a memory operation with a handler.
   /// \param Self is a shared_ptr to this queue.

sycl/source/handler.cpp

@@ -314,26 +314,24 @@ fill_copy_args(detail::handler_impl *impl,
 
 handler::handler(std::shared_ptr<detail::queue_impl> Queue,
                  bool CallerNeedsEvent)
-    : impl(std::make_shared<detail::handler_impl>(Queue.get(), nullptr,
-                                                  CallerNeedsEvent)),
+    : impl(std::make_shared<detail::handler_impl>(nullptr, CallerNeedsEvent)),
       MQueue(std::move(Queue)) {}
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 // TODO: This function is not used anymore, remove it in the next
 // ABI-breaking window.
-handler::handler(std::shared_ptr<detail::queue_impl> Queue,
-                 std::shared_ptr<detail::queue_impl> PrimaryQueue,
-                 std::shared_ptr<detail::queue_impl> SecondaryQueue,
-                 bool CallerNeedsEvent)
-    : impl(std::make_shared<detail::handler_impl>(
-          PrimaryQueue.get(), SecondaryQueue.get(), CallerNeedsEvent)),
+handler::handler(
+    std::shared_ptr<detail::queue_impl> Queue,
+    [[maybe_unused]] std::shared_ptr<detail::queue_impl> PrimaryQueue,
+    std::shared_ptr<detail::queue_impl> SecondaryQueue, bool CallerNeedsEvent)
+    : impl(std::make_shared<detail::handler_impl>(SecondaryQueue.get(),
+                                                  CallerNeedsEvent)),
       MQueue(Queue) {}
 #endif
 
 handler::handler(std::shared_ptr<detail::queue_impl> Queue,
-                 detail::queue_impl *PrimaryQueue,
                  detail::queue_impl *SecondaryQueue, bool CallerNeedsEvent)
-    : impl(std::make_shared<detail::handler_impl>(PrimaryQueue, SecondaryQueue,
+    : impl(std::make_shared<detail::handler_impl>(SecondaryQueue,
                                                   CallerNeedsEvent)),
       MQueue(std::move(Queue)) {}
 
@@ -1769,8 +1767,7 @@ void handler::ext_oneapi_signal_external_semaphore(
 
 void handler::use_kernel_bundle(
     const kernel_bundle<bundle_state::executable> &ExecBundle) {
-  if ((!impl->MGraph && (impl->MSubmissionPrimaryQueue->get_context() !=
-                         ExecBundle.get_context())) ||
+  if ((!impl->MGraph && (MQueue->get_context() != ExecBundle.get_context())) ||
       (impl->MGraph &&
        (impl->MGraph->getContext() != ExecBundle.get_context())))
     throw sycl::exception(
@@ -1931,23 +1928,21 @@ void handler::verifyDeviceHasProgressGuarantee(
 }
 
 bool handler::supportsUSMMemcpy2D() {
-  auto &PrimQueue = impl->MSubmissionPrimaryQueue;
-  if (PrimQueue)
-    return checkContextSupports(PrimQueue->getContextImplPtr(),
-                                UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT);
-  else
-    // Return true when handler_impl is constructed with a graph.
+  // Return true when handler_impl is constructed with a graph.
+  if (!MQueue)
     return true;
+
+  return checkContextSupports(MQueue->getContextImplPtr(),
+                              UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT);
 }
 
 bool handler::supportsUSMFill2D() {
-  auto &PrimQueue = impl->MSubmissionPrimaryQueue;
-  if (PrimQueue)
-    return checkContextSupports(PrimQueue->getContextImplPtr(),
-                                UR_CONTEXT_INFO_USM_FILL2D_SUPPORT);
-  else
-    // Return true when handler_impl is constructed with a graph.
+  // Return true when handler_impl is constructed with a graph.
+  if (!MQueue)
     return true;
+
+  return checkContextSupports(MQueue->getContextImplPtr(),
+                              UR_CONTEXT_INFO_USM_FILL2D_SUPPORT);
 }
 
 bool handler::supportsUSMMemset2D() {

sycl/test/abi/sycl_symbols_windows.dump

@@ -402,7 +402,7 @@
 ??0gpu_selector@_V1@sycl@@QEAA@XZ
 ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z
 ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z
-??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@PEAVqueue_impl@detail@12@1_N@Z
+??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@PEAVqueue_impl@detail@12@_N@Z
 ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z
 ??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z
 ??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z