[WIP] No handler submit POC - header based version

Author: slawekptak

sycl/source/detail/kernel_name_based_cache_t.hpp renamed to sycl/include/sycl/detail/kernel_name_based_cache_t.hpp

@@ -7,8 +7,8 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
-#include <detail/hashers.hpp>
-#include <detail/kernel_arg_mask.hpp>
+//#include <detail/hashers.hpp>
+//#include <sycl/detail/kernel_arg_mask.hpp>
 #include <emhash/hash_table8.hpp>
 #include <sycl/detail/spinlock.hpp>
 #include <sycl/detail/ur.hpp>
@@ -20,6 +20,7 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 using FastKernelCacheKeyT = std::pair<ur_device_handle_t, ur_context_handle_t>;
+using KernelArgMask = std::vector<bool>;
 
 struct FastKernelCacheVal {
   ur_kernel_handle_t MKernelHandle;    /* UR kernel handle pointer. */
@@ -29,25 +30,27 @@ struct FastKernelCacheVal {
   const KernelArgMask *MKernelArgMask; /* Eliminated kernel argument mask. */
   ur_program_handle_t MProgramHandle;  /* UR program handle corresponding to
                                      this kernel. */
-  const Adapter &MAdapterPtr;          /* We can keep reference to the adapter
+  /*const Adapter &MAdapterPtr;*/          /* We can keep reference to the adapter
                                      because during 2-stage shutdown the kernel
                                      cache is destroyed deliberately before the
                                      adapter. */
 
   FastKernelCacheVal(ur_kernel_handle_t KernelHandle, std::mutex *Mutex,
                      const KernelArgMask *KernelArgMask,
-                     ur_program_handle_t ProgramHandle,
-                     const Adapter &AdapterPtr)
+                     ur_program_handle_t ProgramHandle)
+                     //const Adapter &AdapterPtr)
       : MKernelHandle(KernelHandle), MMutex(Mutex),
-        MKernelArgMask(KernelArgMask), MProgramHandle(ProgramHandle),
-        MAdapterPtr(AdapterPtr) {}
+        MKernelArgMask(KernelArgMask), MProgramHandle(ProgramHandle)
+        /*MAdapterPtr(AdapterPtr)*/ {}
 
   ~FastKernelCacheVal() {
+    /*
     if (MKernelHandle)
       MAdapterPtr.call<sycl::detail::UrApiKind::urKernelRelease>(MKernelHandle);
     if (MProgramHandle)
       MAdapterPtr.call<sycl::detail::UrApiKind::urProgramRelease>(
           MProgramHandle);
+    */
     MKernelHandle = nullptr;
     MMutex = nullptr;
     MKernelArgMask = nullptr;

sycl/include/sycl/khr/free_function_commands.hpp

@@ -150,28 +150,28 @@ void launch_grouped(handler &h, range<3> r, range<3> size,
 }
 
 template <typename KernelType>
-void launch_grouped(const queue &q, range<1> r, range<1> size,
+void launch_grouped(queue &q, range<1> r, range<1> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
   (void)codeLoc;
-  q.parallel_for_no_handler(nd_range<1>(r, size), k);
+  q.parallel_for_no_handler_v2(nd_range<1>(r, size), k);
 }
 template <typename KernelType>
-void launch_grouped(const queue &q, range<2> r, range<2> size,
+void launch_grouped(queue &q, range<2> r, range<2> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
   (void)codeLoc;
-  q.parallel_for_no_handler(nd_range<2>(r, size), k);
+  q.parallel_for_no_handler_v2(nd_range<2>(r, size), k);
 }
 template <typename KernelType>
-void launch_grouped(const queue &q, range<3> r, range<3> size,
+void launch_grouped(queue &q, range<3> r, range<3> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
   (void)codeLoc;
-  q.parallel_for_no_handler(nd_range<3>(r, size), k);
+  q.parallel_for_no_handler_v2(nd_range<3>(r, size), k);
 }
 
 template <typename... Args>
@@ -297,6 +297,11 @@ void launch_task(const queue &q, const kernel &k, Args &&...args) {
          [&](handler &h) { launch_task(h, k, std::forward<Args>(args)...); });
 }
 
+template <typename FuncT>
+void launch_host_task(queue &q, FuncT &&Func) {
+  q.host_task_no_handler(std::move(Func));
+}
+
 inline void memcpy(handler &h, void *dest, const void *src, size_t numBytes) {
   h.memcpy(dest, src, numBytes);
 }

sycl/include/sycl/queue.hpp

@@ -40,6 +40,7 @@
 #include <sycl/nd_range.hpp>                         // for nd_range
 #include <sycl/property_list.hpp>                    // for property_list
 #include <sycl/range.hpp>                            // for range
+#include <sycl/detail/kernel_name_based_cache_t.hpp>
 
 #include <cstddef>     // for size_t
 #include <functional>  // for function
@@ -2869,6 +2870,119 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
       KernelNameBasedCachePtr);
   }
 
+  bool check_event_readiness(std::shared_ptr<detail::event_impl> EventImpl) const;
+  ur_device_handle_t get_device_ur_handle() const;
+  ur_context_handle_t get_context_ur_handle() const;
+  void extract_args_set_arg_value(ur_kernel_handle_t Kernel,
+    size_t NextTrueIndex, int Size, const void *ArgPtr) const;
+  void extract_args_set_arg_pointer(ur_kernel_handle_t Kernel,
+    size_t NextTrueIndex, const void *Ptr) const;
+
+  detail::FastKernelCacheValPtr
+  getKernel(ur_device_handle_t Device,
+    ur_context_handle_t Context, detail::FastKernelSubcacheT *KernelSubcacheHint) const {
+
+    const detail::FastKernelSubcacheEntriesT &SubcacheEntries =
+            KernelSubcacheHint->Entries;
+    detail::FastKernelSubcacheReadLockT SubcacheLock{KernelSubcacheHint->Mutex};
+    const detail::FastKernelCacheKeyT RequiredKey(Device, Context);
+    // Search for the kernel in the subcache.
+    auto It = std::find_if(SubcacheEntries.begin(), SubcacheEntries.end(),
+                            [&](const detail::FastKernelEntryT &Entry) {
+                              return Entry.Key == RequiredKey;
+                            });
+    if (It != SubcacheEntries.end()) {
+      //traceKernel("Kernel fetched.", KernelName, true);
+      return It->Value;
+    }
+
+    return detail::FastKernelCacheValPtr();
+  }
+
+  void extract_args_from_lambda(ur_kernel_handle_t Kernel, void *KernelFuncPtr,
+    const detail::KernelArgMask *EliminatedArgMask, int KernelNumParams,
+    detail::kernel_param_desc_t (*KernelParamDescGetter)(int)) const {
+    auto setFunc = [this, Kernel, KernelFuncPtr](const detail::kernel_param_desc_t &ParamDesc,
+                                size_t NextTrueIndex) {
+      const void *ArgPtr = (const char *)KernelFuncPtr + ParamDesc.offset;
+      switch (ParamDesc.kind) {
+      case detail::kernel_param_kind_t::kind_std_layout: {
+        int Size = ParamDesc.info;
+        extract_args_set_arg_value(Kernel, NextTrueIndex, Size, ArgPtr);
+        break;
+      }
+      case detail::kernel_param_kind_t::kind_pointer: {
+        const void *Ptr = *static_cast<const void *const *>(ArgPtr);
+        extract_args_set_arg_pointer(Kernel, NextTrueIndex, Ptr);
+        break;
+      }
+      default:
+        throw std::runtime_error("Direct kernel argument copy failed.");
+      }
+    };
+
+    if (!EliminatedArgMask || EliminatedArgMask->size() == 0) {
+      for (int I = 0; I < KernelNumParams; ++I) {
+        const detail::kernel_param_desc_t &Param = KernelParamDescGetter(I);
+        setFunc(Param, I);
+      }
+    } else {
+      size_t NextTrueIndex = 0;
+      for (int I = 0; I < KernelNumParams; ++I) {
+        const detail::kernel_param_desc_t &Param = KernelParamDescGetter(I);
+        if ((*EliminatedArgMask)[I])
+          continue;
+        setFunc(Param, NextTrueIndex);
+        ++NextTrueIndex;
+      }
+    }
+  }
+
+  template <typename KernelName, typename KernelType, int Dims>
+  void submit_no_handler_v2(nd_range<Dims> Range, const KernelType &KernelFunc) {
+
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+
+    KernelType KernelFuncLocal = KernelFunc;
+    void *KernelFuncPtr = reinterpret_cast<void *>(&KernelFuncLocal);
+    int KernelNumParams = detail::getKernelNumParams<NameT>();
+    detail::kernel_param_desc_t (*KernelParamDescGetter)(int) = &(detail::getKernelParamDesc<NameT>);
+    bool IsKernelESIMD = detail::isKernelESIMD<NameT>();
+    bool HasSpecialCapt = detail::hasSpecialCaptures<NameT>();
+    detail::KernelNameBasedCacheT *KernelNameBasedCachePtr = detail::getKernelNameBasedCache<NameT>();
+    ur_device_handle_t Device = get_device_ur_handle();
+    ur_context_handle_t Context = get_context_ur_handle();
+
+    assert(HasSpecialCapt == false);
+    assert(IsKernelESIMD == false);
+
+    detail::FastKernelCacheValPtr KernelCacheVal = getKernel(
+        Device, Context, &KernelNameBasedCachePtr->FastKernelSubcache);
+
+    std::lock_guard<std::mutex> Lock{*MMutex};
+
+    bool SchedulerBypass = KernelCacheVal &&
+      (!MLastEvent || (check_event_readiness(MLastEvent)));
+
+    if (SchedulerBypass) {
+      extract_args_from_lambda(KernelCacheVal->MKernelHandle, KernelFuncPtr,
+        KernelCacheVal->MKernelArgMask, KernelNumParams, KernelParamDescGetter);
+      submit_no_handler_impl(Range, KernelCacheVal);
+      MLastEvent = nullptr;
+    } else {
+      const char *KernelN = detail::getKernelName<NameT>();
+
+      std::shared_ptr<detail::event_impl> EventImpl = submit_no_handler_impl_v2(
+        Range, KernelN, KernelFuncPtr, KernelNumParams, KernelParamDescGetter,
+        KernelNameBasedCachePtr, MLastEvent);
+
+      if (is_in_order()) {
+        MLastEvent = EventImpl;
+      }
+    }
+  }
+
 public:
   /// single_task version not using handler
   template <typename KernelName = detail::auto_name, typename KernelType>
@@ -2888,6 +3002,29 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     submit_no_handler<KernelName, KernelType, Dims>(Range, KernelFunc);
   }
 
+  template <typename KernelName = detail::auto_name, int Dims,
+            typename KernelType>
+  void parallel_for_no_handler_v2(nd_range<Dims> Range, const KernelType &KernelFunc) {
+
+    kernel_parallel_for<KernelName, sycl::nd_item<Dims>, KernelType,
+      ext::oneapi::experimental::empty_properties_t>(KernelFunc);
+    submit_no_handler_v2<KernelName, KernelType, Dims>(Range, KernelFunc);
+  }
+
+  template <typename FuncT>
+  std::enable_if_t<detail::check_fn_signature<std::remove_reference_t<FuncT>,
+                                              void()>::value ||
+                    detail::check_fn_signature<std::remove_reference_t<FuncT>,
+                                              void(interop_handle)>::value>
+  host_task_no_handler(FuncT &&Func) {
+    std::lock_guard<std::mutex> Lock{*MMutex};
+
+    std::shared_ptr<detail::event_impl> EventImpl =
+      host_task_no_handler_impl(std::move(Func), MLastEvent);
+    if (is_in_order() && EventImpl) {
+      MLastEvent = EventImpl;
+    }
+  }
 
 
   /// parallel_for version with a kernel represented as a lambda + range that
@@ -3673,6 +3810,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::shared_ptr<detail::queue_impl> impl;
   queue(std::shared_ptr<detail::queue_impl> impl) : impl(impl) {}
 
+  std::shared_ptr<detail::event_impl> MLastEvent;
+  std::shared_ptr<std::mutex> MMutex;
+
   template <class Obj>
   friend const decltype(Obj::impl) &
   detail::getSyclObjImpl(const Obj &SyclObject);
@@ -3799,6 +3939,19 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     int KernelNumParams, detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
     detail::KernelNameBasedCacheT *KernelNameBasedCachePtr) const;
 
+  template<int Dims>
+  std::shared_ptr<detail::event_impl> submit_no_handler_impl_v2(
+    nd_range<Dims> Range, const char *KernelName, void *KernelFunc,
+    int KernelNumParams, detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
+    detail::KernelNameBasedCacheT *KernelNameBasedCachePtr,
+    std::shared_ptr<detail::event_impl> LastEvent) const;
+
+  template<int Dims>
+  void submit_no_handler_impl(nd_range<Dims> Range, detail::FastKernelCacheValPtr &KernelCacheVal) const;
+
+  std::shared_ptr<detail::event_impl> host_task_no_handler_impl(
+    std::function<void()> &&Func, std::shared_ptr<detail::event_impl> LastEventImpl);
+
   /// Submits a command group function object to the queue, in order to be
   /// scheduled for execution on the device.
   ///

sycl/source/detail/global_handler.cpp

@@ -14,7 +14,7 @@
 #include <detail/adapter.hpp>
 #include <detail/config.hpp>
 #include <detail/global_handler.hpp>
-#include <detail/kernel_name_based_cache_t.hpp>
+#include <sycl/detail/kernel_name_based_cache_t.hpp>
 #include <detail/platform_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <detail/scheduler/scheduler.hpp>

sycl/source/detail/kernel_program_cache.hpp

@@ -11,7 +11,8 @@
 #include "sycl/exception.hpp"
 #include <detail/config.hpp>
 #include <detail/kernel_arg_mask.hpp>
-#include <detail/kernel_name_based_cache_t.hpp>
+#include <detail/hashers.hpp>
+#include <sycl/detail/kernel_name_based_cache_t.hpp>
 #include <detail/platform_impl.hpp>
 #include <detail/unordered_multimap.hpp>
 #include <sycl/detail/common.hpp>

sycl/source/detail/program_manager/program_manager.cpp

@@ -1186,7 +1186,7 @@ FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
     // nullptr for the mutex.
     auto [Kernel, ArgMask] = BuildF();
     return std::make_shared<FastKernelCacheVal>(
-        Kernel, nullptr, ArgMask, Program, *ContextImpl.getAdapter().get());
+        Kernel, nullptr, ArgMask, Program); //*ContextImpl.getAdapter().get());
   }
 
   auto BuildResult = Cache.getOrBuild<errc::invalid>(GetCachedBuildF, BuildF);
@@ -1195,7 +1195,7 @@ FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
   const KernelArgMaskPairT &KernelArgMaskPair = BuildResult->Val;
   auto ret_val = std::make_shared<FastKernelCacheVal>(
       KernelArgMaskPair.first, &(BuildResult->MBuildResultMutex),
-      KernelArgMaskPair.second, Program, *ContextImpl.getAdapter().get());
+      KernelArgMaskPair.second, Program); //*ContextImpl.getAdapter().get());
   // If caching is enabled, one copy of the kernel handle will be
   // stored in FastKernelCacheVal, and one is in
   // KernelProgramCache::MKernelsPerProgramCache. To cover

sycl/source/detail/program_manager/program_manager.hpp

@@ -12,7 +12,7 @@
 #include <detail/device_global_map_entry.hpp>
 #include <detail/host_pipe_map_entry.hpp>
 #include <detail/kernel_arg_mask.hpp>
-#include <detail/kernel_name_based_cache_t.hpp>
+#include <sycl/detail/kernel_name_based_cache_t.hpp>
 #include <detail/spec_constant_impl.hpp>
 #include <sycl/detail/cg_types.hpp>
 #include <sycl/detail/common.hpp>