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

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

@@ -20,6 +20,10 @@ namespace sycl {
 inline namespace _V1 {
 class handler;
 
+namespace ext::oneapi::experimental::detail {
+class dynamic_parameter_impl;
+}
+
 namespace detail {
 template <typename T> struct is_unbounded_array : std::false_type {};
 
@@ -28,6 +32,18 @@ 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;
+
+void processArg(
+    void *Ptr, const kernel_param_kind_t &Kind, const int Size,
+    const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource,
+    bool IsESIMD, 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 +55,14 @@ class work_group_memory_impl {
 private:
   size_t buffer_size;
   friend class sycl::handler;
+  friend void processArg(
+      void *Ptr, const kernel_param_kind_t &Kind, const int Size,
+      const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource,
+      bool IsESIMD, 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

@@ -501,24 +501,19 @@ class __SYCL_EXPORT handler {
   extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum,
                                const detail::kernel_param_desc_t *KernelArgs,
                                bool IsESIMD);
+  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
+
   /// 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);
-
   /// 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
-  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);
-
   /// \return a string containing name of SYCL kernel.
   detail::ABINeutralKernelNameStrT getKernelName();
 

sycl/include/sycl/stream.hpp

@@ -41,6 +41,15 @@ inline namespace _V1 {
 
 namespace detail {
 
+void processArg(
+    void *Ptr, const kernel_param_kind_t &Kind, const int Size,
+    const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource,
+    bool IsESIMD, 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 +1051,14 @@ 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;