[SYCL] Move Kernel specific data from handler_impl to a separate data structure (#19843)

This PR is a prerequisite for the handler-less API.
Kernel-specific data and argument parsing logic are moved from the
`handler_impl` to the new `KernelData` class that will be used in a
handler-less path.

Author: vinser52

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

@@ -18,9 +18,10 @@
 
 namespace sycl {
 inline namespace _V1 {
-class handler;
 
 namespace detail {
+class KernelData;
+
 template <typename T> struct is_unbounded_array : std::false_type {};
 
 template <typename T> struct is_unbounded_array<T[]> : std::true_type {};
@@ -38,7 +39,7 @@ class work_group_memory_impl {
 
 private:
   size_t buffer_size;
-  friend class sycl::handler;
+  friend class KernelData;
 };
 
 } // namespace detail

sycl/include/sycl/handler.hpp

@@ -495,9 +495,7 @@ class __SYCL_EXPORT handler {
   template <class Kernel> void setDeviceKernelInfo(void *KernelFuncPtr) {
     constexpr auto Info = detail::CompileTimeKernelInfo<Kernel>;
     MKernelName = Info.Name;
-    // TODO support ESIMD in no-integration-header case too.
-    setKernelInfo(KernelFuncPtr, Info.NumParams, Info.ParamDescGetter,
-                  Info.IsESIMD, Info.HasSpecialCaptures);
+    setKernelFunc(KernelFuncPtr);
     setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<Kernel>());
     setType(detail::CGType::Kernel);
   }
@@ -514,23 +512,21 @@ class __SYCL_EXPORT handler {
   extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum,
                                const detail::kernel_param_desc_t *KernelArgs,
                                bool IsESIMD);
-#endif
   /// Extracts and prepares kernel arguments from the lambda using information
   /// from the built-ins or integration header.
   void extractArgsAndReqsFromLambda(
       char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int),
       size_t NumKernelParams, bool IsESIMD);
-
+#endif
   /// Extracts and prepares kernel arguments set via set_arg(s).
   void extractArgsAndReqs();
 
-#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
-  // TODO: processArg need not to be public
-  __SYCL_DLL_LOCAL
-#endif
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // TODO: remove in the next ABI-breaking window.
   void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
                   const int Size, const size_t Index, size_t &IndexShift,
                   bool IsKernelCreatedFromSource, bool IsESIMD);
+#endif
 
   /// \return a string containing name of SYCL kernel.
   detail::ABINeutralKernelNameStrT getKernelName();
@@ -3604,7 +3600,10 @@ class __SYCL_EXPORT handler {
 
   void addArg(detail::kernel_param_kind_t ArgKind, void *Req, int AccessTarget,
               int ArgIndex);
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // TODO: remove in the next ABI-breaking window
   void clearArgs();
+#endif
   void setArgsToAssociatedAccessors();
 
   bool HasAssociatedAccessor(detail::AccessorImplHost *Req,
@@ -3651,10 +3650,12 @@ class __SYCL_EXPORT handler {
   void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset);
   void setNDRangeDescriptor(sycl::range<1> NumWorkItems,
                             sycl::range<1> LocalSize, sycl::id<1> Offset);
-
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
                      detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
                      bool KernelIsESIMD, bool KernelHasSpecialCaptures);
+#endif
+  void setKernelFunc(void *KernelFuncPtr);
 
   void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr);
 

sycl/include/sycl/stream.hpp

@@ -42,6 +42,7 @@ inline namespace _V1 {
 namespace detail {
 
 class stream_impl;
+class KernelData;
 
 using FmtFlags = unsigned int;
 
@@ -1041,7 +1042,7 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream
   }
 #endif
 
-  friend class handler;
+  friend class detail::KernelData;
 
   template <typename SYCLObjT> friend class ext::oneapi::weak_object;
 

sycl/source/CMakeLists.txt

@@ -260,6 +260,7 @@ set(SYCL_COMMON_SOURCES
     "detail/device_filter.cpp"
     "detail/host_pipe_map.cpp"
     "detail/device_global_map.cpp"
+    "detail/kernel_data.cpp"
     "detail/kernel_global_info.cpp"
     "detail/device_global_map_entry.cpp"
     "detail/device_image_impl.cpp"

sycl/source/detail/device_kernel_info.cpp

@@ -63,7 +63,11 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
     const CompileTimeKernelInfoTy &Info) {
   if (!isCompileTimeInfoSet())
     CompileTimeKernelInfoTy::operator=(Info);
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  // In case of 6.3 compatibility mode the KernelSize is not passed to the
+  // runtime. So, it will always be 0 and this assert fails.
   assert(isCompileTimeInfoSet());
+#endif
   assert(Info == *this);
 }
 

sycl/source/detail/graph/dynamic_impl.cpp

@@ -343,7 +343,7 @@ void dynamic_command_group_impl::finalizeCGFList(
     MCommandGroups.push_back(std::shared_ptr<sycl::detail::CG>(RawCGPtr));
 
     // Track dynamic_parameter usage in command-group
-    auto &DynamicParams = Handler.impl->MDynamicParameters;
+    auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters();
 
     if (DynamicParams.size() > 0 &&
         Handler.getType() == sycl::detail::CGType::CodeplayHostTask) {

sycl/source/detail/graph/graph_impl.cpp

@@ -471,7 +471,7 @@ node_impl &graph_impl::add(std::function<void(handler &)> CGF,
 
   // Retrieve any dynamic parameters which have been registered in the CGF and
   // register the actual nodes with them.
-  auto &DynamicParams = Handler.impl->MDynamicParameters;
+  auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters();
 
   if (NodeType != node_type::kernel && DynamicParams.size() > 0) {
     throw sycl::exception(sycl::make_error_code(errc::invalid),

sycl/source/detail/handler_impl.hpp

@@ -11,6 +11,7 @@
 #include "sycl/handler.hpp"
 #include <detail/cg.hpp>
 #include <detail/kernel_bundle_impl.hpp>
+#include <detail/kernel_data.hpp>
 #include <memory>
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
 
@@ -61,8 +62,7 @@ class handler_impl {
   }
 
   KernelNameStrRefT getKernelName() const {
-    assert(MDeviceKernelInfoPtr);
-    return static_cast<KernelNameStrRefT>(MDeviceKernelInfoPtr->Name);
+    return MKernelData.getKernelName();
   }
 
   /// Registers mutually exclusive submission states.
@@ -108,12 +108,6 @@ class handler_impl {
   // If the pipe operation is read or write, 1 for read 0 for write.
   bool HostPipeRead = true;
 
-  ur_kernel_cache_config_t MKernelCacheConfig = UR_KERNEL_CACHE_CONFIG_DEFAULT;
-
-  bool MKernelIsCooperative = false;
-  bool MKernelUsesClusterLaunch = false;
-  uint32_t MKernelWorkGroupMemorySize = 0;
-
   // Extra information for bindless image copy
   ur_image_desc_t MSrcImageDesc = {};
   ur_image_desc_t MDstImageDesc = {};
@@ -138,29 +132,17 @@ class handler_impl {
   sycl::ext::oneapi::experimental::node_type MUserFacingNodeType =
       sycl::ext::oneapi::experimental::node_type::empty;
 
-  // Storage for any SYCL Graph dynamic parameters which have been flagged for
-  // registration in the CG, along with the argument index for the parameter.
-  std::vector<std::pair<
-      ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
-      MDynamicParameters;
-
   /// The storage for the arguments passed.
   /// We need to store a copy of values that are passed explicitly through
   /// set_arg, require and so on, because we need them to be alive after
   /// we exit the method they are passed in.
   detail::CG::StorageInitHelper CGData;
 
-  /// The list of arguments for the kernel.
-  std::vector<detail::ArgDesc> MArgs;
-
   /// The list of associated accessors with this handler.
   /// These accessors were created with this handler as argument or
   /// have become required for this handler via require method.
   std::vector<detail::ArgDesc> MAssociatedAccesors;
 
-  /// Struct that encodes global size, local size, ...
-  detail::NDRDescT MNDRDesc;
-
   /// Type of the command group, e.g. kernel, fill. Can also encode version.
   /// Use getType and setType methods to access this variable unless
   /// manipulations with version are required
@@ -241,16 +223,16 @@ class handler_impl {
   // Allocation ptr to be freed asynchronously.
   void *MFreePtr = nullptr;
 
-  // Store information about the kernel arguments.
-  void *MKernelFuncPtr = nullptr;
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // TODO: remove in the next ABI-breaking window
+  // Today they are used only in the handler::setKernelNameBasedCachePtr
   int MKernelNumArgs = 0;
   detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr;
   bool MKernelIsESIMD = false;
   bool MKernelHasSpecialCaptures = true;
+#endif
 
-  // A pointer to device kernel information. Cached on the application side in
-  // headers or retrieved from program manager.
-  DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr;
+  KernelData MKernelData;
 };
 
 } // namespace detail