Merge branch 'no_handler_single_task' into temp_no_handler_integration_v2

Author: slawekptak

devops/scripts/benchmarks/html/scripts.js

@@ -345,7 +345,7 @@ function createChart(data, containerId, type) {
             if (elements.length > 0) {
                 const point = elements[0].element.$context.raw;
                 if (point.gitHash && point.gitRepo) {
-                    window.open(`https://github.com/${point.gitRepo}/commit/${point.gitHash}`, '_blank');
+                    window.open(`${point.gitRepo}/commit/${point.gitHash}`, '_blank');
                 }
             }
         };

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -170,13 +170,13 @@ class SYCLToolchain {
         llvm::vfs::getRealFileSystem());
     FS->pushOverlay(ToolchainFS);
     if (FSOverlay)
-      FS->pushOverlay(FSOverlay);
+      FS->pushOverlay(std::move(FSOverlay));
 
     auto Files = llvm::makeIntrusiveRefCnt<clang::FileManager>(
         clang::FileSystemOptions{"." /* WorkingDir */}, FS);
 
     Action A{FEAction};
-    ToolInvocation TI{CommandLine, &A, Files.get(),
+    ToolInvocation TI{std::move(CommandLine), &A, Files.get(),
                       std::make_shared<PCHContainerOperations>()};
     TI.setDiagnosticConsumer(DiagConsumer ? DiagConsumer : &IgnoreDiag);
 

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -152,9 +152,13 @@ template <typename KernelName = sycl::detail::auto_name, typename KernelType>
 void single_task(queue Q, const KernelType &KernelObj,
                  const sycl::detail::code_location &CodeLoc =
                      sycl::detail::code_location::current()) {
+  /*
   submit(
       std::move(Q),
       [&](handler &CGH) { single_task<KernelName>(CGH, KernelObj); }, CodeLoc);
+  */
+  detail::submit_kernel_direct_single_task<KernelName>(std::move(Q), empty_properties_t{},
+                                             KernelObj);
 }
 
 template <typename... ArgsT>
@@ -265,7 +269,7 @@ void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                 !(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
                           const KernelType &>::value)) {
-    detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
+    detail::submit_kernel_direct_parallel_for<KernelName>(std::move(Q), empty_properties_t{},
                                              Range, KernelObj);
   } else {
     submit(std::move(Q), [&](handler &CGH) {

sycl/include/sycl/khr/free_function_commands.hpp

@@ -161,7 +161,7 @@ void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
                           const KernelType &>::value)) {
-    detail::submit_kernel_direct(
+    detail::submit_kernel_direct_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<1>(r, size), std::forward<KernelType>(k));
   } else {
@@ -179,7 +179,7 @@ void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
                           const KernelType &>::value)) {
-    detail::submit_kernel_direct(
+    detail::submit_kernel_direct_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<2>(r, size), std::forward<KernelType>(k));
   } else {
@@ -197,7 +197,7 @@ void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
                           const KernelType &>::value)) {
-    detail::submit_kernel_direct(
+    detail::submit_kernel_direct_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<3>(r, size), std::forward<KernelType>(k));
   } else {
@@ -314,7 +314,10 @@ template <typename KernelType>
 void launch_task(const sycl::queue &q, const KernelType &k,
                  const sycl::detail::code_location &codeLoc =
                      sycl::detail::code_location::current()) {
-  submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);
+  //submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);
+  detail::submit_kernel_direct_single_task(q,
+    ext::oneapi::experimental::empty_properties_t{},
+    k, codeLoc);
 }
 
 template <typename... Args>

sycl/include/sycl/queue.hpp

@@ -158,7 +158,7 @@ class __SYCL_EXPORT SubmissionInfo {
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
           typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
-auto submit_kernel_direct(
+auto submit_kernel_direct_parallel_for(
     const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
     KernelTypeUniversalRef &&KernelFunc,
     const detail::code_location &CodeLoc = detail::code_location::current()) {
@@ -211,6 +211,53 @@ auto submit_kernel_direct(
   }
 }
 
+template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+          typename PropertiesT, typename KernelTypeUniversalRef>
+auto submit_kernel_direct_single_task(
+    const queue &Queue, PropertiesT Props, KernelTypeUniversalRef &&KernelFunc,
+    const detail::code_location &CodeLoc = detail::code_location::current()) {
+  // TODO Properties not supported yet
+  (void)Props;
+  static_assert(
+      std::is_same_v<PropertiesT,
+                     ext::oneapi::experimental::empty_properties_t>,
+      "Setting properties not supported yet for no-CGH kernel submit.");
+  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+
+  using KernelType =
+      std::remove_const_t<std::remove_reference_t<KernelTypeUniversalRef>>;
+
+  using NameT =
+      typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+
+  detail::KernelWrapper<detail::WrapAs::single_task, NameT, KernelType,
+                        void, PropertiesT>::wrap(KernelFunc);
+
+  HostKernelRef<KernelType, KernelTypeUniversalRef, void, 1>
+      HostKernel(std::forward<KernelTypeUniversalRef>(KernelFunc));
+
+  // Instantiating the kernel on the host improves debugging.
+  // Passing this pointer to another translation unit prevents optimization.
+#ifndef NDEBUG
+  // TODO: call library to prevent dropping call due to optimization
+  (void)
+      detail::GetInstantiateKernelOnHostPtr<KernelType, void, 1>();
+#endif
+
+  detail::DeviceKernelInfo *DeviceKernelInfoPtr =
+      &detail::getDeviceKernelInfo<NameT>();
+
+  if constexpr (EventNeeded) {
+    return submit_kernel_direct_with_event_impl(
+        Queue, nd_range<1>{1, 1}, HostKernel, DeviceKernelInfoPtr,
+        TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+  } else {
+    submit_kernel_direct_without_event_impl(
+        Queue, nd_range<1>{1, 1}, HostKernel, DeviceKernelInfoPtr,
+        TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+  }
+}
+
 } // namespace detail
 
 namespace ext ::oneapi ::experimental {
@@ -2720,14 +2767,21 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                     void(kernel_handler)>::value),
         "sycl::queue.single_task() requires a kernel instead of command group. "
         "Use queue.submit() instead");
-
+/*
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.template single_task<KernelName, KernelType, PropertiesT>(
               Properties, KernelFunc);
         },
         TlsCodeLocCapture.query());
+*/
+
+      (void)Properties;
+      return detail::submit_kernel_direct_single_task<KernelName, true>(
+          *this, ext::oneapi::experimental::empty_properties_t{},
+          KernelFunc, CodeLoc);
+
   }
 
   /// single_task version with a kernel represented as a lambda.
@@ -3275,7 +3329,6 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
     constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
-
     using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
 
     // TODO The handler-less path does not support reductions and kernel
@@ -3284,7 +3337,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                   !(ext::oneapi::experimental::detail::
                         HasKernelPropertiesGetMethod<
                             const KernelType &>::value)) {
-      return detail::submit_kernel_direct<KernelName, true>(
+      return detail::submit_kernel_direct_parallel_for<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest...);
     } else {

sycl/source/detail/graph/graph_impl.cpp

@@ -549,6 +549,13 @@ graph_impl::add(std::shared_ptr<dynamic_command_group_impl> &DynCGImpl,
   return NodeImpl;
 }
 
+std::shared_ptr<sycl::detail::queue_impl> graph_impl::getQueue() const {
+  std::shared_ptr<sycl::detail::queue_impl> Return{};
+  if (!MRecordingQueues.empty())
+    Return = MRecordingQueues.begin()->lock();
+  return Return;
+}
+
 void graph_impl::addQueue(sycl::detail::queue_impl &RecordingQueue) {
   MRecordingQueues.insert(RecordingQueue.weak_from_this());
 }
@@ -870,10 +877,6 @@ exec_graph_impl::exec_graph_impl(sycl::context Context,
                                  const std::shared_ptr<graph_impl> &GraphImpl,
                                  const property_list &PropList)
     : MSchedule(), MGraphImpl(GraphImpl), MSyncPoints(),
-      MQueueImpl(sycl::detail::queue_impl::create(
-          *sycl::detail::getSyclObjImpl(GraphImpl->getDevice()),
-          *sycl::detail::getSyclObjImpl(Context), sycl::async_handler{},
-          sycl::property_list{})),
       MDevice(GraphImpl->getDevice()), MContext(Context), MRequirements(),
       MSchedulerDependencies(),
       MIsUpdatable(PropList.has_property<property::graph::updatable>()),
@@ -893,6 +896,15 @@ exec_graph_impl::exec_graph_impl(sycl::context Context,
   }
   // Copy nodes from GraphImpl and merge any subgraph nodes into this graph.
   duplicateNodes();
+
+  if (auto PlaceholderQueuePtr = GraphImpl->getQueue()) {
+    MQueueImpl = PlaceholderQueuePtr;
+  } else {
+    MQueueImpl = sycl::detail::queue_impl::create(
+        *sycl::detail::getSyclObjImpl(GraphImpl->getDevice()),
+        *sycl::detail::getSyclObjImpl(Context), sycl::async_handler{},
+        sycl::property_list{});
+  }
 }
 
 exec_graph_impl::~exec_graph_impl() {

sycl/source/detail/graph/graph_impl.hpp

@@ -172,6 +172,8 @@ class graph_impl : public std::enable_shared_from_this<graph_impl> {
   node_impl &add(std::shared_ptr<dynamic_command_group_impl> &DynCGImpl,
                  nodes_range Deps);
 
+  std::shared_ptr<sycl::detail::queue_impl> getQueue() const;
+
   /// Add a queue to the set of queues which are currently recording to this
   /// graph.
   /// @param RecordingQueue Queue to add to set.

sycl/source/detail/queue_impl.cpp

@@ -662,6 +662,14 @@ queue_impl::submit_direct(bool CallerNeedsEvent,
            : true) &&
       !hasCommandGraph();
 
+  if (isInOrder()) {
+    if (SchedulerBypass) {
+      MNoLastEventMode.store(true, std::memory_order_relaxed);
+    } else {
+      MNoLastEventMode.store(false, std::memory_order_relaxed);
+    }
+  }
+
   EventImplPtr EventImpl = SubmitCommandFunc(CGData, SchedulerBypass);
 
   // Sync with the last event for in order queue. For scheduler-bypass flow,

sycl/source/interop_handle.cpp

@@ -60,7 +60,9 @@ ur_native_handle_t interop_handle::getNativeContext() const {
 
 ur_native_handle_t
 interop_handle::getNativeQueue(int32_t &NativeHandleDesc) const {
-  return MQueue->getNative(NativeHandleDesc);
+  if (MQueue != nullptr)
+    return MQueue->getNative(NativeHandleDesc);
+  return 0;
 }
 
 ur_native_handle_t interop_handle::getNativeGraph() const {

sycl/test-e2e/SubGroup/sub_group_as_vec.cpp

@@ -33,7 +33,7 @@ int main(int argc, char *argv[]) {
     queue.submit([&](sycl::handler &cgh) {
       auto global = buf.get_access<sycl::access::mode::read_write,
                                    sycl::access::target::device>(cgh);
-#ifdef DUSE_DEPRECATED_LOCAL_ACC
+#ifdef USE_DEPRECATED_LOCAL_ACC
       sycl::accessor<sycl::vec<int, 2>, 1, sycl::access::mode::read_write,
                      sycl::access::target::local>
           local(N, cgh);