[SYCL][NFC] Add a new header with kernel argument extraction logic

The kernel argument extraction logic has been refactored
and moved to a separate header, to be used in both handler-based
and handler-less kernel submission flows.

Author: slawekptak

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

@@ -28,6 +28,17 @@ template <typename T> struct is_unbounded_array<T[]> : std::true_type {};
 template <typename T>
 inline constexpr bool is_unbounded_array_v = is_unbounded_array<T>::value;
 
+class NDRDescT;
+class ArgDesc;
+class dynamic_parameter_impl;
+
+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,
+                         detail::NDRDescT NDRDesc, std::vector<std::pair<
+                          ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
+                         DynamicParameters, std::vector<ArgDesc> &Args);
+
 class work_group_memory_impl {
 public:
   work_group_memory_impl() : buffer_size{0} {}
@@ -39,6 +50,12 @@ class work_group_memory_impl {
 private:
   size_t buffer_size;
   friend class sycl::handler;
+  friend void detail::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
+                         const int Size, const size_t Index, size_t &IndexShift,
+                         bool IsKernelCreatedFromSource, bool IsESIMD,
+                         detail::NDRDescT NDRDesc, std::vector<std::pair<
+                          ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
+                         DynamicParameters, std::vector<ArgDesc> &Args);
 };
 
 } // namespace detail

sycl/include/sycl/handler.hpp

@@ -499,12 +499,12 @@ 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();

sycl/include/sycl/stream.hpp

@@ -41,6 +41,13 @@ inline namespace _V1 {
 
 namespace detail {
 
+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,
+                         detail::NDRDescT NDRDesc, std::vector<std::pair<
+                          ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
+                         DynamicParameters, std::vector<ArgDesc> &Args);
+
 class stream_impl;
 
 using FmtFlags = unsigned int;
@@ -1042,6 +1049,12 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream
 #endif
 
   friend class handler;
+  friend void detail::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
+                         const int Size, const size_t Index, size_t &IndexShift,
+                         bool IsKernelCreatedFromSource, bool IsESIMD,
+                         detail::NDRDescT NDRDesc, std::vector<std::pair<
+                          ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
+                         DynamicParameters, std::vector<ArgDesc> &Args);
 
   template <typename SYCLObjT> friend class ext::oneapi::weak_object;