From 15555b7ea783b1648531aa931d0e568881b3f6f3 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Fri, 18 Jul 2025 19:24:45 +0200
Subject: [PATCH 01/23] [SYCL] add possibility to get number of argumnet of
 kernell free function

---
 clang/lib/Sema/SemaSYCL.cpp                   |  32 +-
 ...ee_function_default_template_arguments.cpp | 461 ++++++++++++-
 .../CodeGenSYCL/free_function_int_header.cpp  | 605 ++++++++++++++++--
 .../free_function_int_header_rtc_mode.cpp     |   5 +-
 sycl/include/sycl/detail/kernel_desc.hpp      |   5 +
 .../sycl/ext/oneapi/get_kernel_info.hpp       |   9 +
 sycl/include/sycl/kernel.hpp                  |   4 +
 sycl/include/sycl/kernel_bundle.hpp           |   6 +-
 sycl/source/detail/kernel_impl.cpp            |  16 +
 sycl/source/detail/kernel_impl.hpp            |   9 +-
 sycl/source/kernel.cpp                        |   4 +
 .../test-e2e/FreeFunctionKernels/num_args.cpp |  68 ++
 sycl/test/abi/sycl_symbols_linux.dump         |   1 +
 13 files changed, 1137 insertions(+), 88 deletions(-)
 create mode 100644 sycl/test-e2e/FreeFunctionKernels/num_args.cpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3feebc423e2aa..8329b419665c3 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6666,6 +6666,32 @@ class FreeFunctionPrinter {
           FD->getTemplateSpecializationArgs());
   }
 
+  /// Emits free function kernel info specialization for shimN.
+  /// \param ShimCounter The counter for the shim function.
+  /// \param KParamsSize The number of kernel free function arguments.
+  /// \param KName The name of the kernel free function.
+  void printFreeFunctionKernelInfo(const unsigned ShimCounter,
+                                   const size_t KParamsSize,
+                                   std::string_view KName) {
+    O << "\nnamespace sycl {\n";
+    O << "inline namespace _V1 {\n";
+    O << "namespace detail {\n";
+    O << "//Free Function Kernel info specialization for shim" << ShimCounter
+      << "\n";
+    O << "template <> struct FreeFunctionInfoData<__sycl_shim" << ShimCounter
+      << "()> {\n";
+    O << "\t__SYCL_DLL_LOCAL\n";
+    O << "\tstatic constexpr unsigned getNumParams() { return " << KParamsSize
+      << "; }\n";
+    O << "\t__SYCL_DLL_LOCAL\n";
+    O << "\tstatic constexpr const char *getFunctionName() { return ";
+    O << "\"" << KName << "\"; }\n";
+    O << "};\n";
+    O << "} // namespace detail\n"
+      << "} // namespace _V1\n"
+      << "} // namespace sycl\n\n";
+  }
+
 private:
   /// Helper method to get string with template types
   /// \param TAL The template argument list.
@@ -7127,6 +7153,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     FFPrinter.printFreeFunctionShim(K.SyclKernel, ShimCounter, ParmList);
     O << ";\n";
     O << "}\n";
+    FFPrinter.printFreeFunctionKernelInfo(ShimCounter, K.Params.size(), K.Name);
     Policy.SuppressDefaultTemplateArgs = true;
     Policy.EnforceDefaultTemplateArgs = false;
 
@@ -7167,8 +7194,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "template <>\n";
     O << "inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim"
       << ShimCounter << "()>() {\n";
-    O << "  return sycl::detail::get_kernel_id_impl(std::string_view{\""
-      << K.Name << "\"});\n";
+    O << "  return sycl::detail::get_kernel_id_impl(std::string_view{"
+      << "sycl::detail::FreeFunctionInfoData<__sycl_shim" << ShimCounter
+      << "()>::getFunctionName()});\n";
     O << "}\n";
     O << "}\n";
     ++ShimCounter;
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 2debb64fcc8a3..45ba5c8026f07 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -256,6 +256,20 @@ namespace Testing::Tests {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<char, int, 12, struct ns::notatuple>))ns::simple;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim1
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim1()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT:	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT:	__SYCL_DLL_LOCAL
+// CHECK-NEXT:	static constexpr const char *getFunctionName() { return "_ZN16__sycl_kernel_ns6simpleENS_3ArgIciLi12ENS_9notatupleEJEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: Forward declarations of kernel and its argument types:
 // CHECK: namespace ns {
 // CHECK: namespace ns1 {
@@ -267,21 +281,96 @@ namespace Testing::Tests {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, 12, struct ns::notatuple>))simple1;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim2
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim2()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z21__sycl_kernel_simple1N2ns3ArgINS_3ns113hasDefaultArgINS_9notatupleEEEiLi12ES3_JEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple> , T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim3() {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<int, float, 3, struct ns::notatuple>, int))templated<int>;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim3
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim3()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z23__sycl_kernel_templatedIiEvN2ns3ArgIT_fLi3ENS0_9notatupleEJEEES2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: template <typename T> void templated2(ns::Arg<T, ns::notatuple, 12, ns::notatuple> , T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim4() {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<int, struct ns::notatuple, 12, struct ns::notatuple>, int))templated2<int>;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim4
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim4()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z24__sycl_kernel_templated2IiEvN2ns3ArgIT_NS0_9notatupleELi12ES3_JEEES2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: template <typename T, int a> void templated3(ns::Arg<T, ns::notatuple, a, ns::ns1::hasDefaultArg<ns::notatuple>, int, int> , T end);
 // CHECK-NEXT: static constexpr auto __sycl_shim5() {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<int, struct ns::notatuple, 3, class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, int>, int))templated3<int, 3>;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim5
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim5()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z24__sycl_kernel_templated3IiLi3EEvN2ns3ArgIT_NS0_9notatupleEXT0_ENS0_3ns113hasDefaultArgIS3_EEJiiEEES2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: template <typename T, int a> void templated3(ns::Arg<T, ns::notatuple, a, ns::ns1::hasDefaultArg<ns::notatuple>, int, int> , T end);
+// CHECK-NEXT: static constexpr auto __sycl_shim6() {
+// CHECK-NEXT:   return (void (*)(struct ns::Arg<float, struct ns::notatuple, 3, class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, int>, float))templated3<float, 3>;
+// CHECK-NEXT: }
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim6
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim6()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z24__sycl_kernel_templated3IfLi3EEvN2ns3ArgIT_NS0_9notatupleEXT0_ENS0_3ns113hasDefaultArgIS3_EEJiiEEES2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK Forward declarations of kernel and its argument types:
 // CHECK: namespace sycl { namespace detail {
 // CHECK-NEXT: struct Y;
@@ -296,6 +385,20 @@ namespace Testing::Tests {
 // CHECK-NEXT:   return (void (*)(struct Arg1<int, struct sycl::X<struct sycl::detail::Y> >))foo;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim7
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim7()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z17__sycl_kernel_foo4Arg1IiN4sycl1XINS0_6detail1YEEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace TestNamespace {
 // CHECK-NEXT:  template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple> , T end);
 // CHECK-NEXT:  } // namespace TestNamespace
@@ -303,7 +406,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim8() {
 // CHECK-NEXT:    return (void (*)(struct ns::Arg<int, float, 3, struct ns::notatuple>, int))TestNamespace::templated<int>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim8
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim8()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN27__sycl_kernel_TestNamespace9templatedIiEEvN2ns3ArgIT_fLi3ENS1_9notatupleEJEEES3_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -322,7 +440,22 @@ namespace Testing::Tests {
 // CHECK:   static constexpr auto __sycl_shim9() {
 // CHECK-NEXT:    return (void (*)(struct ns::Arg<int, float, 10, struct ns::notatuple>, int))TestNamespace::_V1::templated1<int, 10>;
 // CHECK-NEXT:   }
-// CHECK-NEXT:   namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim9
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim9()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN27__sycl_kernel_TestNamespace3_V110templated1IiLi10EEEvN2ns3ArgIT_fXT0_ENS2_9notatupleEJEEES4_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:   namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim9()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -341,7 +474,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim10() {
 // CHECK-NEXT:    return (void (*)(struct ns::Arg<int, int, 12, struct ns::notatuple>, int))TestNamespace::_V2::templated1<int, 12>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim10
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim10()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN27__sycl_kernel_TestNamespace3_V210templated1IiLi12EEEvN2ns3ArgIT_S4_XT0_ENS2_9notatupleEJEEES4_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -358,7 +506,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim11() {
 // CHECK-NEXT:    return (void (*)(float, float))templated<float>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim11
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim11()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedIfEEvT_S1_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim11()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -374,7 +537,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim12() {
 // CHECK-NEXT:  return (void (*)(struct ns::Arg<struct TestStruct, float, 3, struct ns::notatuple>, struct TestStruct))templated<struct TestStruct>;
 // CHECK-NEXT:}
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim12
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim12()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z23__sycl_kernel_templatedI10TestStructEvN2ns3ArgIT_fLi3ENS1_9notatupleEJEEES3_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim12()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -392,7 +570,22 @@ namespace Testing::Tests {
 // CHECK: static constexpr auto __sycl_shim13() {
 // CHECK-NEXT:  return (void (*)(class BaseClass, class BaseClass))templated<class BaseClass>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim13
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim13()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedI9BaseClassEEvT_S2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim13()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -410,7 +603,22 @@ namespace Testing::Tests {
 // CHECK: static constexpr auto __sycl_shim14() {
 // CHECK-NEXT:  return (void (*)(class ChildOne, class ChildOne))templated<class ChildOne>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim14
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim14()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedI8ChildOneEEvT_S2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim14()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -428,7 +636,22 @@ namespace Testing::Tests {
 // CHECK: static constexpr auto __sycl_shim15() {
 // CHECK-NEXT:  return (void (*)(class ChildTwo, class ChildTwo))templated<class ChildTwo>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim15
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim15()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedI8ChildTwoEEvT_S2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim15()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -446,7 +669,22 @@ namespace Testing::Tests {
 // CHECK: static constexpr auto __sycl_shim16() {
 // CHECK-NEXT:  return (void (*)(class ChildThree, class ChildThree))templated<class ChildThree>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim16
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim16()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedI10ChildThreeEEvT_S2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim16()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -466,7 +704,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim17() {
 // CHECK-NEXT:    return (void (*)(struct sycl::id<2>, struct sycl::id<2>))templated<struct sycl::id<2>>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim17
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim17()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedIN4sycl3_V12idILi2EEEEEvT_S5_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim17()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -486,7 +739,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim18() {
 // CHECK-NEXT:    return (void (*)(struct sycl::range<3>, struct sycl::range<3>))templated<struct sycl::range<3>>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim18
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim18()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedIN4sycl3_V15rangeILi3EEEEEvT_S5_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim18()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -503,7 +771,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim19() {
 // CHECK-NEXT:    return (void (*)(int *, int *))templated<int *>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim19
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim19()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedIPiEEvT_S2_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim19()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -520,7 +803,23 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim20() {
 // CHECK-NEXT:    return (void (*)(struct sycl::X<class ChildTwo>, struct sycl::X<class ChildTwo>))templated<struct sycl::X<class ChildTwo>>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim20
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim20()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_19templatedIN4sycl1XI8ChildTwoEEEEvT_S5_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim20()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -542,7 +841,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim21() {
 // CHECK-NEXT:    return (void (*)(struct ns::Arg<struct One::Two::Three::AnotherStruct, float, 10, struct ns::notatuple>, struct One::Two::Three::AnotherStruct))TestNamespace::_V1::templated1<struct One::Two::Three::AnotherStruct, 10>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim21
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim21()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN27__sycl_kernel_TestNamespace3_V110templated1IN3One3Two5Three13AnotherStructELi10EEEvN2ns3ArgIT_fXT0_ENS6_9notatupleEJEEES8_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim21()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -557,7 +871,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim22() {
 // CHECK-NEXT:  return (void (*)(int, float, char))variadic_templated<int, float, char>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim22
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim22()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z32__sycl_kernel_variadic_templatedIJifcEEvDpT_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim22()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -572,7 +901,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim23() {
 // CHECK-NEXT:  return (void (*)(int, float, char, int))variadic_templated<int, float, char, int>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim23
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim23()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 4; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z32__sycl_kernel_variadic_templatedIJifciEEvDpT_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim23()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -587,7 +931,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim24() {
 // CHECK-NEXT:  return (void (*)(float, float))variadic_templated<float, float>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim24
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim24()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z32__sycl_kernel_variadic_templatedIJffEEvDpT_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim24()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -602,7 +961,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim25() {
 // CHECK-NEXT:  return (void (*)(float, char, char))variadic_templated1<float, char, char>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim25
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim25()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z33__sycl_kernel_variadic_templated1IfJccEEvT_DpT0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim25()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -617,7 +991,22 @@ namespace Testing::Tests {
 // CHECK-NEXT: static constexpr auto __sycl_shim26() {
 // CHECK-NEXT:  return (void (*)(int, float, char))variadic_templated1<int, float, char>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim26
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim26()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z33__sycl_kernel_variadic_templated1IiJfcEEvT_DpT0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim26()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -636,7 +1025,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim27() {
 // CHECK-NEXT:    return (void (*)(float, float))Testing::Tests::variadic_templated<float, float>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim27
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim27()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN21__sycl_kernel_Testing5Tests18variadic_templatedIfJfEEEvT_DpT0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim27()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -655,7 +1059,22 @@ namespace Testing::Tests {
 // CHECK:  static constexpr auto __sycl_shim28() {
 // CHECK-NEXT:    return (void (*)(int, int, int, int))Testing::Tests::variadic_templated<int, int, int, int>;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim28
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim28()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 4; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN21__sycl_kernel_Testing5Tests18variadic_templatedIiJiiiEEEvT_DpT0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim28()> {
 // CHECK-NEXT:    static constexpr bool value = true;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index 48a03c6c65916..34ebc6668558b 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -448,7 +448,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim1() {
 // CHECK-NEXT:   return (void (*)(int *, int, int))ff_2;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim1
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim1()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_2Piii"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim1()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -465,7 +480,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim2() {
 // CHECK-NEXT:   return (void (*)(int *, int, int, int))ff_2;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim2
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim2()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 4; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_2Piiii"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim2()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -482,7 +512,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim3() {
 // CHECK-NEXT:   return (void (*)(int *, int, int))ff_3<int>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim3
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim3()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_3IiEvPT_S0_S0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim3()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -499,7 +544,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim4() {
 // CHECK-NEXT:   return (void (*)(float *, float, float))ff_3<float>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim4
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim4()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_3IfEvPT_S0_S0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim4()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -517,7 +577,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim5() {
 // CHECK-NEXT:   return (void (*)(double *, double, double))ff_3<double>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim5
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim5()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim5()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -537,7 +612,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim6() {
 // CHECK-NEXT:   return (void (*)(struct NoPointers, struct Pointers, struct Agg))ff_4;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim6
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim6()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim6()> {
 // CHECK-NEXT:   static constexpr bool value = true;
@@ -555,7 +645,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim7() {
 // CHECK-NEXT:   return (void (*)(struct Agg, struct Derived, int))ff_6<struct Agg, struct Derived>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim7
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim7()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 3; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim7()> {
 // CHECK-NEXT:   static constexpr bool value = true;
@@ -575,7 +680,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim8() {
 // CHECK-NEXT:   return (void (*)(struct KArgWithPtrArray<3>))ff_7<3>;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim8
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim8()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
 // CHECK-NEXT:   static constexpr bool value = true;
@@ -595,7 +715,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim9() {
 // CHECK-NEXT: return (void (*)(class sycl::work_group_memory<int>))ff_8;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim9
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim9()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim9()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -617,6 +752,20 @@ void ff_24(int arg) {
 // CHECK-NEXT:   return (void (*)(int, int *))free_functions::ff_9;
 // CHECK-NEXT: }
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim10
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim10()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions4ff_9EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi as a free function kernel
 // CHECK: Forward declarations of kernel and its argument types:
 
@@ -628,7 +777,22 @@ void ff_24(int arg) {
 // CHECK: static constexpr auto __sycl_shim11() {
 // CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_10;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim11
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim11()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5tests5ff_10EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim11()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -652,7 +816,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim12() {
 // CHECK-NEXT:    return (void (*)(int, int *))free_functions::tests::V1::ff_11;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim12
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim12()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim12()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -672,7 +851,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim13() {
 // CHECK-NEXT:    return (void (*)(int, int *))ff_12;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim13
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim13()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim13()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -692,7 +886,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim14() {
 // CHECK-NEXT:    return (void (*)(int, int *))free_functions::ff_13;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim14
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim14()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5ff_13EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim14()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -714,7 +923,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim15() {
 // CHECK-NEXT:    return (void (*)(int, int *))free_functions::tests::ff_13;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim15
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim15()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5tests5ff_13EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim15()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -736,8 +960,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim16() {
 // CHECK-NEXT: return (void (*)(class sycl::dynamic_work_group_memory<int>))ff_9;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim16
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim16()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim16()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -753,9 +991,23 @@ void ff_24(int arg) {
 // CHECK-NEXT: template <typename dataT, int dimensions> class local_accessor;
 
 // CHECK: void ff_11(sycl::local_accessor<int, 1> lacc);
-// CHECK-NEXT: static constexpr auto __sycl_shim
+// CHECK-NEXT: static constexpr auto __sycl_shim17
 // CHECK-NEXT:  return (void (*)(class sycl::local_accessor<int, 1>))ff_11;
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim17
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim17()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_11N4sycl3_V114local_accessorIiLi1EEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -773,6 +1025,20 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim
 // CHECK-NEXT:  return (void (*)(class sycl::local_accessor<float, 1>))ff_11<float>;
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim18
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim18()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_11IfEvN4sycl3_V114local_accessorIT_Li1EEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -792,6 +1058,20 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim
 // CHECK-NEXT: return (void (*)(class sycl::sampler))ff_12;
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim19
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim19()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_12N4sycl3_V17samplerE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -811,6 +1091,20 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim
 // CHECK-NEXT:  return (void (*)(class sycl::stream))ff_13;
 
+// CHECK: namespace sycl {
+// CHECK: inline namespace _V1 {
+// CHECK: namespace detail {
+// CHECK: //Free Function Kernel info specialization for shim20
+// CHECK: template <> struct FreeFunctionInfoData<__sycl_shim20()> {
+// CHECK: 	__SYCL_DLL_LOCAL
+// CHECK: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK: 	__SYCL_DLL_LOCAL
+// CHECK: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_13N4sycl3_V16streamE"; }
+// CHECK: };
+// CHECK: } // namespace detail
+// CHECK: } // namespace _V1
+// CHECK: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -830,6 +1124,20 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim
 // CHECK-NEXT:  return (void (*)(class sycl::ext::oneapi::experimental::annotated_arg<int>))ff_14;
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim21
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim21()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_14N4sycl3_V13ext6oneapi12experimental13annotated_argIiJEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -849,6 +1157,20 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim
 // CHECK-NEXT:  return (void (*)(class sycl::ext::oneapi::experimental::annotated_ptr<int>))ff_15;
 
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim22
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim22()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_15N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel
@@ -869,7 +1191,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim23() {
 // CHECK-NEXT:    return (void (*)(int, int *))free_functions::tests::ff_14;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim23
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim23()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5tests5ff_14EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim23()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -888,7 +1225,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim24() {
 // CHECK-NEXT:    return (void (*)(int, int *))free_functions::ff_15;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim24
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim24()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5ff_15EiPi"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim24()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -907,7 +1259,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim25() {
 // CHECK-NEXT:    return (void (*)(struct Agg, struct Agg *))free_functions::ff_16;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim25
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim25()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim25()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -926,7 +1293,22 @@ void ff_24(int arg) {
 // CHECK:  static constexpr auto __sycl_shim26() {
 // CHECK-NEXT:    return (void (*)(struct Derived, struct Derived *))free_functions::ff_17;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim26
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim26()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim26()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -950,7 +1332,22 @@ void ff_24(int arg) {
 // CHECK:   static constexpr auto __sycl_shim27() {
 // CHECK-NEXT:     return (void (*)(struct free_functions::Agg, struct free_functions::Agg *))free_functions::tests::ff_18;
 // CHECK-NEXT:   }
-// CHECK-NEXT:   namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim27
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim27()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:   namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim27()> {
 // CHECK-NEXT:     static constexpr bool value = true;
@@ -971,7 +1368,22 @@ void ff_24(int arg) {
 // CHECK-NEXT:  static constexpr auto __sycl_shim28() {
 // CHECK-NEXT:    return (void (*)(struct free_functions::KArgWithPtrArray<50>))ff_19;
 // CHECK-NEXT:  }
-// CHECK-NEXT:  namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim28
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim28()> {
+// CHECK-NEXT: __SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim28()> {
 // CHECK-NEXT:    static constexpr bool value = true;
@@ -1004,7 +1416,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim29() {
 // CHECK-NEXT:  return (void (*)(class sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t, class sycl::ext::oneapi::accessor_property_list<> >))ff_20;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim29
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim29()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim29()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -1018,7 +1445,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim30() {
 // CHECK-NEXT:  return (void (*)(struct Derived, struct Derived *))ff_21;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim30
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim30()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_217DerivedPS_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim30()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -1033,7 +1475,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim31() {
 // CHECK-NEXT:  return (void (*)(struct Derived, struct Derived *))ff_22;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim31
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim31()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 2; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_227DerivedPS_"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim31()> {
 // CHECK-NEXT:  static constexpr bool value = true;
@@ -1049,7 +1506,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim32() {
 // CHECK-NEXT:  return (void (*)(int))ff_24;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim32
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim32()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_24i"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim32()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -1065,7 +1537,22 @@ void ff_24(int arg) {
 // CHECK-NEXT: static constexpr auto __sycl_shim33() {
 // CHECK-NEXT: return (void (*)(int))ff_23;
 // CHECK-NEXT: }
-// CHECK-NEXT: namespace sycl {
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: //Free Function Kernel info specialization for shim33
+// CHECK-NEXT: template <> struct FreeFunctionInfoData<__sycl_shim33()> {
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr unsigned getNumParams() { return 1; }
+// CHECK-NEXT: 	__SYCL_DLL_LOCAL
+// CHECK-NEXT: 	static constexpr const char *getFunctionName() { return "_Z19__sycl_kernel_ff_23i"; }
+// CHECK-NEXT: };
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
+
+// CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim33()> {
 // CHECK-NEXT: static constexpr bool value = true;
@@ -1083,7 +1570,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim1()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_2Piii"});
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim1()>::getFunctionName()});
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
@@ -1091,7 +1578,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim2()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_2Piiii"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
@@ -1099,7 +1586,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim3()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IiEvPT_S0_S0_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim3()>::getFunctionName()});
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
@@ -1107,7 +1594,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim4()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IfEvPT_S0_S0_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim4()>::getFunctionName()});
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
@@ -1115,7 +1602,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:   template <>
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim5()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim5()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1123,7 +1610,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim6()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim6()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1131,7 +1618,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim7()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim7()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1139,7 +1626,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim8()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1147,7 +1634,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim9()>() {
-// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim9()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1156,7 +1643,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions4ff_9EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim10()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1164,7 +1651,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim11()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_10EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim11()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  } 
 
@@ -1172,7 +1659,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim12()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi"});
+// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim12()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1180,7 +1667,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim13()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim13()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1188,7 +1675,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim14()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_13EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim14()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1196,7 +1683,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim15()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_13EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim15()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1205,43 +1692,43 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim16()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim16()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_11N4sycl3_V114local_accessorIiLi1EEE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim17()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_11N4sycl3_V114local_accessorIiLi1EEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim17()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_11IfEvN4sycl3_V114local_accessorIT_Li1EEE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim18()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_11IfEvN4sycl3_V114local_accessorIT_Li1EEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim18()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_12N4sycl3_V17samplerE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim19()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_12N4sycl3_V17samplerE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim19()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_13N4sycl3_V16streamE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim20()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_13N4sycl3_V16streamE"});
+// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim20()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_14N4sycl3_V13ext6oneapi12experimental13annotated_argIiJEEE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim21()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_14N4sycl3_V13ext6oneapi12experimental13annotated_argIiJEEE"})
+// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim21()>::getFunctionName()});
 
 // CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_15N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim22()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_15N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim22()>::getFunctionName()});
 
 // CHECK-NEXT: }
 // CHECK-NEXT: }
@@ -1250,7 +1737,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim23()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_14EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim23()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
   
@@ -1258,7 +1745,7 @@ void ff_24(int arg) {
 // CHECK-NEXT:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim24()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_15EiPi"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim24()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
   
@@ -1266,7 +1753,7 @@ void ff_24(int arg) {
 // CHECK-NEXT:  namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim25()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim25()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1274,7 +1761,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim26()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim26()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1282,7 +1769,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT:  template <>
 // CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim27()>() {
-// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim27()>::getFunctionName()});
 // CHECK-NEXT:  }
 // CHECK-NEXT:  }
 
@@ -1290,27 +1777,27 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim29()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim29()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim30()>() {
-// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_217DerivedPS_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim30()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
 // CHECK: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim31()>() {
-// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_227DerivedPS_"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim31()>::getFunctionName()});
 
 // CHECK: // Definition of kernel_id of _Z19__sycl_kernel_ff_24i
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim32()>() {
-// CHECK-NEXT    return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_24i"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim32()>::getFunctionName()});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
@@ -1318,7 +1805,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: template <>
 // CHECK-NEXT: inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim33()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_23i"});
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim33()>::getFunctionName()});
 
 // CHECK-NEXT: }
 // CHECK-NEXT: }
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 29b697691f445..4586c4c806666 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -74,7 +74,6 @@ int main(){
 // CHECK-NORTC-NEXT:  static constexpr bool value = true;
 
 // CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() {
-// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_singlePiii"});
-
+// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim1()>::getFunctionName()});
 // CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() {
-// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_nd_rangePiii"});
+// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});
\ No newline at end of file
diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp
index a3324cb567d25..09d294d1b2d9e 100644
--- a/sycl/include/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/sycl/detail/kernel_desc.hpp
@@ -174,6 +174,11 @@ template <class KernelNameType> struct KernelInfo {
 };
 #endif //__SYCL_UNNAMED_LAMBDA__
 
+template <auto *Func> struct FreeFunctionInfoData {
+  static constexpr unsigned getNumParams() { return 0; }
+  static constexpr const char *getFunctionName() { return ""; }
+};
+
 // Built-ins accept an object due to lacking infrastructure support for
 // accepting types. The kernel name type itself isn't used because it might be
 // incomplete, cv-qualified, or not default constructible. Passing an object
diff --git a/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
index 5f7a97736aee2..9fb570db9670a 100644
--- a/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
+++ b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
@@ -85,6 +85,15 @@ std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>,
 get_kernel_info(const queue &q) {
   return get_kernel_info<Func, Param>(q.get_context(), q.get_device());
 }
+
+template <auto *Func, typename Param>
+std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func> &&
+                     std::is_same_v<Param, sycl::info::kernel::num_args>,
+                 typename sycl::detail::is_kernel_info_desc<Param>::return_type>
+get_kernel_info(const context &, const device &) {
+  return sycl::detail::FreeFunctionInfoData<Func>::getNumParams();
+}
+
 } // namespace experimental
 } // namespace ext::oneapi
 } // namespace _V1
diff --git a/sycl/include/sycl/kernel.hpp b/sycl/include/sycl/kernel.hpp
index b332c09e4772d..b396d13c8aa8b 100644
--- a/sycl/include/sycl/kernel.hpp
+++ b/sycl/include/sycl/kernel.hpp
@@ -249,6 +249,10 @@ class __SYCL_EXPORT kernel : public detail::OwnerLessBase<kernel> {
   typename detail::is_kernel_queue_specific_info_desc<Param>::return_type
   ext_oneapi_get_info(queue Queue, const range<1> &WG) const;
 
+  /// Set kernel free function argument number.
+  /// \param Num is the number of arguments of the free function kernel.
+  void setFreeFuncKernelArgNum(unsigned Num);
+
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.
   kernel(std::shared_ptr<detail::kernel_impl> Impl);
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index aa6cb7bc1d161..8d08b3b1c929c 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -508,8 +508,12 @@ class kernel_bundle : public detail::kernel_bundle_plain,
             typename = std::enable_if_t<_State == bundle_state::executable>>
   std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>, kernel>
   ext_oneapi_get_kernel() {
-    return detail::kernel_bundle_plain::get_kernel(
+    sycl::kernel Kernel = detail::kernel_bundle_plain::get_kernel(
         ext::oneapi::experimental::get_kernel_id<Func>());
+    const unsigned FreeFuncKernelArgNum =
+        sycl::detail::FreeFunctionInfoData<Func>::getNumParams();
+    Kernel.setFreeFuncKernelArgNum(FreeFuncKernelArgNum);
+    return Kernel;
   }
 
   /////////////////////////
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 0cb679f1f0fc3..eaa35522ec0c8 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -113,6 +113,19 @@ bool kernel_impl::isBuiltInKernel(device_impl &Device) const {
       [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
 }
 
+bool kernel_impl::isFreeFunctionKernel() const {
+  const auto ids = MKernelBundleImpl->get_kernel_ids();
+  return std::any_of(ids.begin(), ids.end(), [this](const kernel_id &Id) {
+    const std::string KernelName = Id.get_name();
+    const auto pos = KernelName.find("__sycl_kernel_");
+    return pos != std::string::npos;
+  });
+}
+
+void kernel_impl::setKerenlFreeFuncArgNum(unsigned Num) {
+  FreeFuncKernelArgNum = Num;
+}
+
 void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
   if (isInteropOrSourceBased())
     return;
@@ -122,6 +135,9 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       }))
     return;
 
+  if (isFreeFunctionKernel())
+    return;
+
   throw sycl::exception(
       sycl::make_error_code(errc::invalid),
       "info::kernel::num_args descriptor may only be used to query a kernel "
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index 0c3c1ab0bf1e4..d109f06b572b4 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -242,6 +242,8 @@ class kernel_impl {
   std::mutex *getCacheMutex() const { return MCacheMutex; }
   std::string_view getName() const;
 
+  void setKerenlFreeFuncArgNum(unsigned Num);
+
 private:
   ur_kernel_handle_t MKernel = nullptr;
   const std::shared_ptr<context_impl> MContext;
@@ -254,8 +256,10 @@ class kernel_impl {
   const KernelArgMask *MKernelArgMaskPtr;
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
+  unsigned FreeFuncKernelArgNum = 0;
 
   bool isBuiltInKernel(device_impl &Device) const;
+  bool isFreeFunctionKernel() const;
   void checkIfValidForNumArgsInfoQuery() const;
 
   /// Check if the occupancy limits are exceeded for the given kernel launch
@@ -309,9 +313,10 @@ template <typename Param>
 inline typename Param::return_type kernel_impl::get_info() const {
   static_assert(is_kernel_info_desc<Param>::value,
                 "Invalid kernel information descriptor");
-  if constexpr (std::is_same_v<Param, info::kernel::num_args>)
+  if constexpr (std::is_same_v<Param, info::kernel::num_args>) {
     checkIfValidForNumArgsInfoQuery();
-
+    return FreeFuncKernelArgNum;
+  }
   return get_kernel_info<Param>(this->getHandleRef(), getAdapter());
 }
 
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index 9b566a566c353..fc2fc6d7596f8 100644
--- a/sycl/source/kernel.cpp
+++ b/sycl/source/kernel.cpp
@@ -50,6 +50,10 @@ kernel::get_kernel_bundle() const {
       kernel_bundle<sycl::bundle_state::executable>>(impl->get_kernel_bundle());
 }
 
+void kernel::setFreeFuncKernelArgNum(unsigned Num) {
+  impl->setKerenlFreeFuncArgNum(Num);
+}
+
 template <typename Param>
 detail::ABINeutralT_t<typename detail::is_kernel_info_desc<Param>::return_type>
 kernel::get_info_impl() const {
diff --git a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
new file mode 100644
index 0000000000000..68e57fe534c8e
--- /dev/null
+++ b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
@@ -0,0 +1,68 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %{build} %level_zero_options -o %t.ze.out
+// RUN: %{run} %t.ze.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+static constexpr auto FFTestMark = "Free function Kernel Test:";
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void kernel_func(sycl::item<1> idx, float value, sycl::accessor<int, 1> acc) {}
+
+template <auto *Func>
+int test_num_args_free_function_api(sycl::context &ctxt, sycl::device &dev,
+                                    const int expected_num_args) {
+  const int actual =
+      syclexp::get_kernel_info<Func, sycl::info::kernel::num_args>(ctxt, dev);
+  const bool res = actual == expected_num_args;
+  if (!res)
+    std::cout << FFTestMark << "test_num_args failed: expected_num_args "
+              << expected_num_args << "actual " << actual << std::endl;
+  return res ? 0 : 1;
+}
+
+template <auto *Func>
+int test_num_args_kernel_api(sycl::context &ctxt, sycl::device &dev,
+                             const int expected_num_args) {
+  auto bundle =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  const int actual = bundle.template ext_oneapi_get_kernel<Func>()
+                         .template get_info<sycl::info::kernel::num_args>();
+  std::cout << FFTestMark << "actual number of args: " << actual
+            << " expected: " << expected_num_args << std::endl;
+  const bool res = actual == expected_num_args;
+  if (!res)
+    std::cout << FFTestMark
+              << "test_num_args_kernel_api failed: expected_num_args "
+              << expected_num_args << "actual " << actual << std::endl;
+  return res ? 0 : 1;
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctx = q.get_context();
+  sycl::device dev = q.get_device();
+
+  int ret = test_num_args_free_function_api<func_range>(ctx, dev, 2);
+  ret |= test_num_args_free_function_api<func_single>(ctx, dev, 2);
+  ret |= test_num_args_free_function_api<kernel_func>(ctx, dev, 3);
+  ret |= test_num_args_kernel_api<func_range>(ctx, dev, 2);
+  ret |= test_num_args_kernel_api<func_single>(ctx, dev, 2);
+  ret |= test_num_args_kernel_api<kernel_func>(ctx, dev, 3);
+  return ret;
+}
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index a617484c11041..25b57caf56eb3 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3393,6 +3393,7 @@ _ZN4sycl3_V16kernelC1EP10_cl_kernelRKNS0_7contextE
 _ZN4sycl3_V16kernelC1ESt10shared_ptrINS0_6detail11kernel_implEE
 _ZN4sycl3_V16kernelC2EP10_cl_kernelRKNS0_7contextE
 _ZN4sycl3_V16kernelC2ESt10shared_ptrINS0_6detail11kernel_implEE
+_ZN4sycl3_V16kernel23setFreeFuncKernelArgNumEj
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE

From e954028ce3071893d72292ba3adf923d5fa60ab2 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Fri, 18 Jul 2025 19:27:18 +0200
Subject: [PATCH 02/23] [SYCL] update formatting

---
 clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 4586c4c806666..fd0016846e568 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -76,4 +76,4 @@ int main(){
 // CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() {
 // CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim1()>::getFunctionName()});
 // CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() {
-// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});
\ No newline at end of file
+// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});

From ecfea215ae8895bd48da71a5c0fa2278ec750b4d Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Fri, 18 Jul 2025 19:41:15 +0200
Subject: [PATCH 03/23] [SYCL] do no capture unused parameters in lambda

---
 sycl/source/detail/kernel_impl.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index eaa35522ec0c8..5fe7a26973160 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -115,7 +115,7 @@ bool kernel_impl::isBuiltInKernel(device_impl &Device) const {
 
 bool kernel_impl::isFreeFunctionKernel() const {
   const auto ids = MKernelBundleImpl->get_kernel_ids();
-  return std::any_of(ids.begin(), ids.end(), [this](const kernel_id &Id) {
+  return std::any_of(ids.begin(), ids.end(), [](const kernel_id &Id) {
     const std::string KernelName = Id.get_name();
     const auto pos = KernelName.find("__sycl_kernel_");
     return pos != std::string::npos;

From c0f2b29f31866a3cf3544a38b426ba8d0aa71340 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Fri, 18 Jul 2025 20:20:57 +0200
Subject: [PATCH 04/23] [SYCL][E2E] update win abi test

---
 sycl/test/abi/sycl_symbols_windows.dump | 1 +
 1 file changed, 1 insertion(+)

diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index e87045fc6ae1f..a913e1b183f3c 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4036,6 +4036,7 @@
 ?get@context@_V1@sycl@@QEBAPEAU_cl_context@@XZ
 ?get@device@_V1@sycl@@QEBAPEAU_cl_device_id@@XZ
 ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ
+?setFreeFuncKernelArgNum@kernel@_V1@sycl@@QEAAXI@Z
 ?get@platform@_V1@sycl@@QEBAPEAU_cl_platform_id@@XZ
 ?get@queue@_V1@sycl@@QEBAPEAU_cl_command_queue@@XZ
 ?getAccData@AccessorBaseHost@detail@_V1@sycl@@QEAAAEAUAccHostDataT@234@XZ

From cdc214c7e4625f3103b3c4cdfb473de9273400ba Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Mon, 21 Jul 2025 13:43:13 +0200
Subject: [PATCH 05/23] [SYCL] return free function kernel data only if kernel
 is free function

---
 sycl/include/sycl/kernel.hpp        |  4 ++++
 sycl/include/sycl/kernel_bundle.hpp |  1 +
 sycl/source/detail/kernel_impl.cpp  | 15 ++++++++-------
 sycl/source/detail/kernel_impl.hpp  |  6 ++++--
 sycl/source/kernel.cpp              |  4 ++++
 5 files changed, 21 insertions(+), 9 deletions(-)

diff --git a/sycl/include/sycl/kernel.hpp b/sycl/include/sycl/kernel.hpp
index b396d13c8aa8b..3ce35f0a330af 100644
--- a/sycl/include/sycl/kernel.hpp
+++ b/sycl/include/sycl/kernel.hpp
@@ -253,6 +253,10 @@ class __SYCL_EXPORT kernel : public detail::OwnerLessBase<kernel> {
   /// \param Num is the number of arguments of the free function kernel.
   void setFreeFuncKernelArgNum(unsigned Num);
 
+  /// Update free function kernel cache to return the correct values of info
+  /// requests
+  void updateFreeFuncKernelCache();
+
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.
   kernel(std::shared_ptr<detail::kernel_impl> Impl);
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 8d08b3b1c929c..98d201b08368b 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -512,6 +512,7 @@ class kernel_bundle : public detail::kernel_bundle_plain,
         ext::oneapi::experimental::get_kernel_id<Func>());
     const unsigned FreeFuncKernelArgNum =
         sycl::detail::FreeFunctionInfoData<Func>::getNumParams();
+    Kernel.updateFreeFuncKernelCache();
     Kernel.setFreeFuncKernelArgNum(FreeFuncKernelArgNum);
     return Kernel;
   }
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 5fe7a26973160..070713234137e 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -113,13 +113,14 @@ bool kernel_impl::isBuiltInKernel(device_impl &Device) const {
       [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
 }
 
-bool kernel_impl::isFreeFunctionKernel() const {
+void kernel_impl::updateFreeFuncKernelCache() {
   const auto ids = MKernelBundleImpl->get_kernel_ids();
-  return std::any_of(ids.begin(), ids.end(), [](const kernel_id &Id) {
-    const std::string KernelName = Id.get_name();
-    const auto pos = KernelName.find("__sycl_kernel_");
-    return pos != std::string::npos;
-  });
+  isFreeFuncKernel =
+      std::any_of(ids.begin(), ids.end(), [](const kernel_id &Id) {
+        const std::string KernelName = Id.get_name();
+        const auto pos = KernelName.find("__sycl_kernel_");
+        return pos != std::string::npos;
+      });
 }
 
 void kernel_impl::setKerenlFreeFuncArgNum(unsigned Num) {
@@ -135,7 +136,7 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       }))
     return;
 
-  if (isFreeFunctionKernel())
+  if (isFreeFuncKernel)
     return;
 
   throw sycl::exception(
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index d109f06b572b4..16dcc27b043db 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -243,6 +243,7 @@ class kernel_impl {
   std::string_view getName() const;
 
   void setKerenlFreeFuncArgNum(unsigned Num);
+  void updateFreeFuncKernelCache();
 
 private:
   ur_kernel_handle_t MKernel = nullptr;
@@ -257,9 +258,9 @@ class kernel_impl {
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
   unsigned FreeFuncKernelArgNum = 0;
+  bool isFreeFuncKernel = false;
 
   bool isBuiltInKernel(device_impl &Device) const;
-  bool isFreeFunctionKernel() const;
   void checkIfValidForNumArgsInfoQuery() const;
 
   /// Check if the occupancy limits are exceeded for the given kernel launch
@@ -315,7 +316,8 @@ inline typename Param::return_type kernel_impl::get_info() const {
                 "Invalid kernel information descriptor");
   if constexpr (std::is_same_v<Param, info::kernel::num_args>) {
     checkIfValidForNumArgsInfoQuery();
-    return FreeFuncKernelArgNum;
+    if (isFreeFuncKernel)
+      return FreeFuncKernelArgNum;
   }
   return get_kernel_info<Param>(this->getHandleRef(), getAdapter());
 }
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index fc2fc6d7596f8..930e82f8f4115 100644
--- a/sycl/source/kernel.cpp
+++ b/sycl/source/kernel.cpp
@@ -54,6 +54,10 @@ void kernel::setFreeFuncKernelArgNum(unsigned Num) {
   impl->setKerenlFreeFuncArgNum(Num);
 }
 
+void kernel::updateFreeFuncKernelCache() {
+  impl->updateFreeFuncKernelCache();
+}
+
 template <typename Param>
 detail::ABINeutralT_t<typename detail::is_kernel_info_desc<Param>::return_type>
 kernel::get_info_impl() const {

From a3e94507910b42572444af35d1fd193835bf0573 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Mon, 21 Jul 2025 13:57:23 +0200
Subject: [PATCH 06/23] [SYCL] fix formatting

---
 sycl/source/kernel.cpp | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index 930e82f8f4115..52402185722ad 100644
--- a/sycl/source/kernel.cpp
+++ b/sycl/source/kernel.cpp
@@ -54,9 +54,7 @@ void kernel::setFreeFuncKernelArgNum(unsigned Num) {
   impl->setKerenlFreeFuncArgNum(Num);
 }
 
-void kernel::updateFreeFuncKernelCache() {
-  impl->updateFreeFuncKernelCache();
-}
+void kernel::updateFreeFuncKernelCache() { impl->updateFreeFuncKernelCache(); }
 
 template <typename Param>
 detail::ABINeutralT_t<typename detail::is_kernel_info_desc<Param>::return_type>

From 237344663afc4bdb434948f8449443e4d83fd23a Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Mon, 21 Jul 2025 14:45:29 +0200
Subject: [PATCH 07/23] [SYCL][TEST] update ABI tests

---
 sycl/test/abi/sycl_symbols_linux.dump   | 1 +
 sycl/test/abi/sycl_symbols_windows.dump | 1 +
 2 files changed, 2 insertions(+)

diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 25b57caf56eb3..270acd2f0d2ca 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3394,6 +3394,7 @@ _ZN4sycl3_V16kernelC1ESt10shared_ptrINS0_6detail11kernel_implEE
 _ZN4sycl3_V16kernelC2EP10_cl_kernelRKNS0_7contextE
 _ZN4sycl3_V16kernelC2ESt10shared_ptrINS0_6detail11kernel_implEE
 _ZN4sycl3_V16kernel23setFreeFuncKernelArgNumEj
+_ZN4sycl3_V16kernel25updateFreeFuncKernelCacheEv
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index a913e1b183f3c..d1dbb6784a798 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4037,6 +4037,7 @@
 ?get@device@_V1@sycl@@QEBAPEAU_cl_device_id@@XZ
 ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ
 ?setFreeFuncKernelArgNum@kernel@_V1@sycl@@QEAAXI@Z
+?updateFreeFuncKernelCache@kernel@_V1@sycl@@QEAAXXZ
 ?get@platform@_V1@sycl@@QEBAPEAU_cl_platform_id@@XZ
 ?get@queue@_V1@sycl@@QEBAPEAU_cl_command_queue@@XZ
 ?getAccData@AccessorBaseHost@detail@_V1@sycl@@QEAAAEAUAccHostDataT@234@XZ

From aab52c6c4f92e09bab48106f50df80304a2e126d Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Tue, 22 Jul 2025 10:22:31 +0200
Subject: [PATCH 08/23] [SYCL] do not use public interface to get num_args

---
 clang/lib/Sema/SemaSYCL.cpp               | 14 ++++++------
 sycl/include/sycl/kernel.hpp              |  8 -------
 sycl/include/sycl/kernel_bundle.hpp       | 10 ++++-----
 sycl/source/detail/kernel_bundle_impl.hpp | 16 ++++++++++++++
 sycl/source/detail/kernel_impl.cpp        | 26 +++++++++++------------
 sycl/source/detail/kernel_impl.hpp        |  7 ++----
 sycl/source/kernel.cpp                    |  6 ------
 sycl/source/kernel_bundle.cpp             |  5 +++++
 sycl/test/abi/sycl_symbols_linux.dump     |  3 +--
 9 files changed, 49 insertions(+), 46 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8329b419665c3..ec5847619370d 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6673,23 +6673,25 @@ class FreeFunctionPrinter {
   void printFreeFunctionKernelInfo(const unsigned ShimCounter,
                                    const size_t KParamsSize,
                                    std::string_view KName) {
-    O << "\nnamespace sycl {\n";
+    O << "\n";
+    O << "namespace sycl {\n";
     O << "inline namespace _V1 {\n";
     O << "namespace detail {\n";
     O << "//Free Function Kernel info specialization for shim" << ShimCounter
       << "\n";
     O << "template <> struct FreeFunctionInfoData<__sycl_shim" << ShimCounter
       << "()> {\n";
-    O << "\t__SYCL_DLL_LOCAL\n";
-    O << "\tstatic constexpr unsigned getNumParams() { return " << KParamsSize
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr unsigned getNumParams() { return " << KParamsSize
       << "; }\n";
-    O << "\t__SYCL_DLL_LOCAL\n";
-    O << "\tstatic constexpr const char *getFunctionName() { return ";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr const char *getFunctionName() { return ";
     O << "\"" << KName << "\"; }\n";
     O << "};\n";
     O << "} // namespace detail\n"
       << "} // namespace _V1\n"
-      << "} // namespace sycl\n\n";
+      << "} // namespace sycl\n";
+    O << "\n";
   }
 
 private:
diff --git a/sycl/include/sycl/kernel.hpp b/sycl/include/sycl/kernel.hpp
index 3ce35f0a330af..b332c09e4772d 100644
--- a/sycl/include/sycl/kernel.hpp
+++ b/sycl/include/sycl/kernel.hpp
@@ -249,14 +249,6 @@ class __SYCL_EXPORT kernel : public detail::OwnerLessBase<kernel> {
   typename detail::is_kernel_queue_specific_info_desc<Param>::return_type
   ext_oneapi_get_info(queue Queue, const range<1> &WG) const;
 
-  /// Set kernel free function argument number.
-  /// \param Num is the number of arguments of the free function kernel.
-  void setFreeFuncKernelArgNum(unsigned Num);
-
-  /// Update free function kernel cache to return the correct values of info
-  /// requests
-  void updateFreeFuncKernelCache();
-
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.
   kernel(std::shared_ptr<detail::kernel_impl> Impl);
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 98d201b08368b..534f749471ca8 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -292,6 +292,8 @@ class __SYCL_EXPORT kernel_bundle_plain {
   // a value different from default value.
   bool is_specialization_constant_set(const char *SpecName) const noexcept;
 
+  void addFreeFuncKernelArgsSize(unsigned Size, const kernel_id &KernelID);
+
   detail::KernelBundleImplPtr impl;
 
 private:
@@ -508,13 +510,11 @@ class kernel_bundle : public detail::kernel_bundle_plain,
             typename = std::enable_if_t<_State == bundle_state::executable>>
   std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>, kernel>
   ext_oneapi_get_kernel() {
-    sycl::kernel Kernel = detail::kernel_bundle_plain::get_kernel(
-        ext::oneapi::experimental::get_kernel_id<Func>());
+    auto KernelID = ext::oneapi::experimental::get_kernel_id<Func>();
     const unsigned FreeFuncKernelArgNum =
         sycl::detail::FreeFunctionInfoData<Func>::getNumParams();
-    Kernel.updateFreeFuncKernelCache();
-    Kernel.setFreeFuncKernelArgNum(FreeFuncKernelArgNum);
-    return Kernel;
+    addFreeFuncKernelArgsSize(FreeFuncKernelArgNum, KernelID);
+    return detail::kernel_bundle_plain::get_kernel(KernelID);
   }
 
   /////////////////////////
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 6ce4d38a3420a..339ba0b4c77d7 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -28,6 +28,7 @@
 #include <cstring>
 #include <memory>
 #include <unordered_set>
+#include <unordered_map>
 #include <vector>
 
 #include "split_string.hpp"
@@ -1047,6 +1048,20 @@ class kernel_bundle_impl
 
   DeviceGlobalMap &getDeviceGlobalMap() { return MDeviceGlobals; }
 
+  void AddKernelArgsSize(const std::string &KernelName, unsigned Size) {
+    auto It = MFreeFuncKernelArgsSizeMap.find(KernelName);
+    if (It == MFreeFuncKernelArgsSizeMap.end()) {
+      MFreeFuncKernelArgsSizeMap[KernelName] = Size;
+    }
+  }
+
+  unsigned GetKernelArgsSize(const std::string &KernelName) const {
+    auto It = MFreeFuncKernelArgsSizeMap.find(KernelName);
+    if (It == MFreeFuncKernelArgsSizeMap.end())
+      return 0;
+    return It->second;
+  }
+
 private:
   DeviceGlobalMapEntry *getDeviceGlobalEntry(const std::string &Name) const {
     if (!hasSourceBasedImages() && !hasSYCLBINImages()) {
@@ -1109,6 +1124,7 @@ class kernel_bundle_impl
 
   context MContext;
   std::vector<device_impl *> MDevices;
+  std::unordered_map<std::string, unsigned> MFreeFuncKernelArgsSizeMap;
 
   // For sycl_jit, building from source may have produced sycl binaries that
   // the kernel_bundles now manage.
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 070713234137e..bccac77366679 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -37,6 +37,9 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &Context,
 
   // Enable USM indirect access for interoperability kernels.
   enableUSMIndirectAccess();
+  const std::string KernelName = get_info<info::kernel::function_name>();
+  const auto pos = KernelName.find("__sycl_kernel_");
+  isFreeFuncKernel = pos != std::string::npos;
 }
 
 kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
@@ -56,6 +59,10 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
   // path.
   if (MCreatedFromSource || MIsInterop)
     enableUSMIndirectAccess();
+
+  const std::string KernelName = get_info<info::kernel::function_name>();
+  const auto pos = KernelName.find("__sycl_kernel_");
+  isFreeFuncKernel = pos != std::string::npos;
 }
 
 kernel_impl::~kernel_impl() {
@@ -113,20 +120,6 @@ bool kernel_impl::isBuiltInKernel(device_impl &Device) const {
       [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
 }
 
-void kernel_impl::updateFreeFuncKernelCache() {
-  const auto ids = MKernelBundleImpl->get_kernel_ids();
-  isFreeFuncKernel =
-      std::any_of(ids.begin(), ids.end(), [](const kernel_id &Id) {
-        const std::string KernelName = Id.get_name();
-        const auto pos = KernelName.find("__sycl_kernel_");
-        return pos != std::string::npos;
-      });
-}
-
-void kernel_impl::setKerenlFreeFuncArgNum(unsigned Num) {
-  FreeFuncKernelArgNum = Num;
-}
-
 void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
   if (isInteropOrSourceBased())
     return;
@@ -146,6 +139,11 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       "interoperability function or to query a device built-in kernel");
 }
 
+unsigned kernel_impl ::getFreeFuncKernelArgSize() const {
+  const std::string KernelName = get_info<info::kernel::function_name>();
+  return MKernelBundleImpl->GetKernelArgsSize(KernelName);
+}
+
 void kernel_impl::enableUSMIndirectAccess() const {
   if (!MContext->getPlatformImpl().supports_usm())
     return;
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index 16dcc27b043db..6ee9e74d57c28 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -242,9 +242,6 @@ class kernel_impl {
   std::mutex *getCacheMutex() const { return MCacheMutex; }
   std::string_view getName() const;
 
-  void setKerenlFreeFuncArgNum(unsigned Num);
-  void updateFreeFuncKernelCache();
-
 private:
   ur_kernel_handle_t MKernel = nullptr;
   const std::shared_ptr<context_impl> MContext;
@@ -257,7 +254,6 @@ class kernel_impl {
   const KernelArgMask *MKernelArgMaskPtr;
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
-  unsigned FreeFuncKernelArgNum = 0;
   bool isFreeFuncKernel = false;
 
   bool isBuiltInKernel(device_impl &Device) const;
@@ -275,6 +271,7 @@ class kernel_impl {
                                size_t DynamicLocalMemorySize) const;
 
   void enableUSMIndirectAccess() const;
+  unsigned getFreeFuncKernelArgSize() const;
 };
 
 template <int Dimensions>
@@ -317,7 +314,7 @@ inline typename Param::return_type kernel_impl::get_info() const {
   if constexpr (std::is_same_v<Param, info::kernel::num_args>) {
     checkIfValidForNumArgsInfoQuery();
     if (isFreeFuncKernel)
-      return FreeFuncKernelArgNum;
+      return getFreeFuncKernelArgSize();
   }
   return get_kernel_info<Param>(this->getHandleRef(), getAdapter());
 }
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index 52402185722ad..9b566a566c353 100644
--- a/sycl/source/kernel.cpp
+++ b/sycl/source/kernel.cpp
@@ -50,12 +50,6 @@ kernel::get_kernel_bundle() const {
       kernel_bundle<sycl::bundle_state::executable>>(impl->get_kernel_bundle());
 }
 
-void kernel::setFreeFuncKernelArgNum(unsigned Num) {
-  impl->setKerenlFreeFuncArgNum(Num);
-}
-
-void kernel::updateFreeFuncKernelCache() { impl->updateFreeFuncKernelCache(); }
-
 template <typename Param>
 detail::ABINeutralT_t<typename detail::is_kernel_info_desc<Param>::return_type>
 kernel::get_info_impl() const {
diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
index 3476c8747102f..96448366b0330 100644
--- a/sycl/source/kernel_bundle.cpp
+++ b/sycl/source/kernel_bundle.cpp
@@ -162,6 +162,11 @@ size_t kernel_bundle_plain::ext_oneapi_get_device_global_size(
       std::string(std::string_view(name)));
 }
 
+void kernel_bundle_plain::addFreeFuncKernelArgsSize(unsigned Size,
+                                              const kernel_id &KernelID) {                          
+  impl->AddKernelArgsSize(KernelID.get_name(), Size);
+}
+
 //////////////////////////////////
 ///// sycl::detail free functions
 //////////////////////////////////
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 270acd2f0d2ca..4708e32ca4df1 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3393,8 +3393,6 @@ _ZN4sycl3_V16kernelC1EP10_cl_kernelRKNS0_7contextE
 _ZN4sycl3_V16kernelC1ESt10shared_ptrINS0_6detail11kernel_implEE
 _ZN4sycl3_V16kernelC2EP10_cl_kernelRKNS0_7contextE
 _ZN4sycl3_V16kernelC2ESt10shared_ptrINS0_6detail11kernel_implEE
-_ZN4sycl3_V16kernel23setFreeFuncKernelArgNumEj
-_ZN4sycl3_V16kernel25updateFreeFuncKernelCacheEv
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_5queueENS0_3usm5allocERKNS0_6detail13code_locationE
 _ZN4sycl3_V16mallocEmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
@@ -3844,6 +3842,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain33contains_specialization_constantsEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain3endEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5beginEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
+_ZN4sycl3_V16detail19kernel_bundle_plain25addFreeFuncKernelArgsSizeEjRKNS0_9kernel_idE
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv

From e415843efe362dc9fe89e1bc0a10e9ebee539578 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Tue, 22 Jul 2025 10:35:16 +0200
Subject: [PATCH 09/23] [SYCL] remove win symbols from abi test

---
 sycl/source/kernel_bundle.cpp           | 2 +-
 sycl/test/abi/sycl_symbols_windows.dump | 2 --
 2 files changed, 1 insertion(+), 3 deletions(-)

diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
index 96448366b0330..dc1fdd5ea6790 100644
--- a/sycl/source/kernel_bundle.cpp
+++ b/sycl/source/kernel_bundle.cpp
@@ -163,7 +163,7 @@ size_t kernel_bundle_plain::ext_oneapi_get_device_global_size(
 }
 
 void kernel_bundle_plain::addFreeFuncKernelArgsSize(unsigned Size,
-                                              const kernel_id &KernelID) {                          
+                                                    const kernel_id &KernelID) {
   impl->AddKernelArgsSize(KernelID.get_name(), Size);
 }
 
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index d1dbb6784a798..e87045fc6ae1f 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4036,8 +4036,6 @@
 ?get@context@_V1@sycl@@QEBAPEAU_cl_context@@XZ
 ?get@device@_V1@sycl@@QEBAPEAU_cl_device_id@@XZ
 ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ
-?setFreeFuncKernelArgNum@kernel@_V1@sycl@@QEAAXI@Z
-?updateFreeFuncKernelCache@kernel@_V1@sycl@@QEAAXXZ
 ?get@platform@_V1@sycl@@QEBAPEAU_cl_platform_id@@XZ
 ?get@queue@_V1@sycl@@QEBAPEAU_cl_command_queue@@XZ
 ?getAccData@AccessorBaseHost@detail@_V1@sycl@@QEAAAEAUAccHostDataT@234@XZ

From 9e1f16d4738be31cc66f7b3216a16524deab7744 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Tue, 22 Jul 2025 10:41:12 +0200
Subject: [PATCH 10/23] [SYCL] fix syntax

---
 sycl/source/detail/kernel_bundle_impl.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 339ba0b4c77d7..de5daeda0c0ab 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -27,8 +27,8 @@
 #include <cstdint>
 #include <cstring>
 #include <memory>
-#include <unordered_set>
 #include <unordered_map>
+#include <unordered_set>
 #include <vector>
 
 #include "split_string.hpp"

From c0fdbc42887ffe533de1540628afdca98fa9796c Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 12:34:11 +0200
Subject: [PATCH 11/23] [SYCL] store kernel global data in a separate map in
 program manager

---
 clang/lib/Sema/SemaSYCL.cpp                   | 37 +++++++++++++++++++
 .../int-header-empty-signatures.cpp           |  2 +
 clang/test/CodeGenSYCL/integration_header.cpp |  2 +
 .../CodeGenSYCL/kernel-param-acc-array-ih.cpp |  2 +
 .../kernel-param-member-acc-array-ih.cpp      |  2 +
 .../CodeGenSYCL/kernel-param-pod-array-ih.cpp |  2 +
 .../CodeGenSYCL/union-kernel-param-ih.cpp     |  2 +
 clang/test/CodeGenSYCL/wrapped-accessor.cpp   |  2 +
 .../sycl/detail/kernel_global_info.hpp        | 23 ++++++++++++
 sycl/include/sycl/kernel_bundle.hpp           | 10 +----
 sycl/source/CMakeLists.txt                    |  1 +
 sycl/source/detail/kernel_bundle_impl.hpp     | 19 ++++------
 sycl/source/detail/kernel_global_info.cpp     | 22 +++++++++++
 .../program_manager/program_manager.cpp       | 13 +++++++
 .../program_manager/program_manager.hpp       | 14 ++++++-
 sycl/source/kernel_bundle.cpp                 |  5 ---
 .../test-e2e/FreeFunctionKernels/num_args.cpp | 37 ++++++++++++++-----
 sycl/test/abi/sycl_symbols_linux.dump         |  2 +-
 18 files changed, 161 insertions(+), 36 deletions(-)
 create mode 100644 sycl/include/sycl/detail/kernel_global_info.hpp
 create mode 100644 sycl/source/detail/kernel_global_info.cpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ec5847619370d..91d7e54570889 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6943,6 +6943,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "  \"\",\n";
   O << "};\n\n";
 
+  O << "static constexpr unsigned kernel_args_sizes[] = {";
+  for (unsigned I = 0; I < KernelDescs.size(); I++) {
+    O << KernelDescs[I].Params.size() << ", ";
+  }
+  O << "};\n\n";
   O << "// array representing signatures of all kernels defined in the\n";
   O << "// corresponding source\n";
   O << "static constexpr\n";
@@ -7203,6 +7208,38 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "}\n";
     ++ShimCounter;
   }
+  O << "#include <sycl/detail/kernel_global_info.hpp>\n";
+
+  ShimCounter = 0;
+  O << "namespace sycl {\n";
+  O << "inline namespace _V1 {\n";
+  O << "namespace detail {\n";
+  O << "namespace free_function_map {\n";
+  O << "inline void update_device_global_map() {\n";
+  for (const KernelDesc &K : KernelDescs) {
+    if (!S.isFreeFunction(K.SyclKernel))
+      continue;
+    O << "sycl::detail::free_function_info_map::add("
+      << "reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + "
+      << ShimCounter << "), sycl::detail::kernel_names[" << ShimCounter
+      << "]);\n";
+    ++ShimCounter;
+  }
+  O << "}\n";
+  O << "struct GlobalMapUpdater {\n";
+  O << "  GlobalMapUpdater() {\n";
+  O << "    update_device_global_map();\n";
+  O << "    std::cout << \"Device global map updated for free function "
+       "kernels. "
+       "Total: "
+    << ShimCounter << " kernels.\" << std::endl;\n";
+  O << "  }\n";
+  O << "};\n";
+  O << "static GlobalMapUpdater updater;\n";
+  O << "} // namespace free_function_map\n";
+  O << "} // namespace detail\n";
+  O << "} // namespace _V1\n";
+  O << "} // namespace sycl\n";
 }
 
 bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) {
diff --git a/clang/test/CodeGenSYCL/int-header-empty-signatures.cpp b/clang/test/CodeGenSYCL/int-header-empty-signatures.cpp
index 88a1c7b182c85..d1b54648f2645 100644
--- a/clang/test/CodeGenSYCL/int-header-empty-signatures.cpp
+++ b/clang/test/CodeGenSYCL/int-header-empty-signatures.cpp
@@ -10,6 +10,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE1K
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index 624bf52dba085..df700e1b66c02 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -25,6 +25,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 //
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+//
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE12first_kernel
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
index 90e24fb115918..c5948f48519f1 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
@@ -18,6 +18,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE8kernel_A
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
index a71d2d9abc28a..a7a115f394be8 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
@@ -18,6 +18,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
index b06bb58c1b28a..1c55478122b2c 100644
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
@@ -19,6 +19,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
diff --git a/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp b/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp
index 4f219c386b528..ab9359a95ce6f 100644
--- a/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp
+++ b/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp
@@ -18,6 +18,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE8kernel_A
diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
index c47faec09f134..a45d285ba4ce5 100644
--- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp
+++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
@@ -15,6 +15,8 @@
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
+// CHECK: static constexpr unsigned kernel_args_sizes[] = {
+
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE14wrapped_access
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
new file mode 100644
index 0000000000000..a6b6f1593f0b9
--- /dev/null
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -0,0 +1,23 @@
+//==-------------------- kernel_global_info.hpp -----------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/export.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+namespace free_function_info_map {
+
+__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
+
+} // namespace free_function_info_map
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
\ No newline at end of file
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 534f749471ca8..fbe7024d412a1 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -291,9 +291,6 @@ class __SYCL_EXPORT kernel_bundle_plain {
   // \returns a bool value which indicates if specialization constant was set to
   // a value different from default value.
   bool is_specialization_constant_set(const char *SpecName) const noexcept;
-
-  void addFreeFuncKernelArgsSize(unsigned Size, const kernel_id &KernelID);
-
   detail::KernelBundleImplPtr impl;
 
 private:
@@ -510,11 +507,8 @@ class kernel_bundle : public detail::kernel_bundle_plain,
             typename = std::enable_if_t<_State == bundle_state::executable>>
   std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>, kernel>
   ext_oneapi_get_kernel() {
-    auto KernelID = ext::oneapi::experimental::get_kernel_id<Func>();
-    const unsigned FreeFuncKernelArgNum =
-        sycl::detail::FreeFunctionInfoData<Func>::getNumParams();
-    addFreeFuncKernelArgsSize(FreeFuncKernelArgNum, KernelID);
-    return detail::kernel_bundle_plain::get_kernel(KernelID);
+    return detail::kernel_bundle_plain::get_kernel(
+        ext::oneapi::experimental::get_kernel_id<Func>());
   }
 
   /////////////////////////
diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
index dd1b833383055..3b9126e4a3bfa 100644
--- a/sycl/source/CMakeLists.txt
+++ b/sycl/source/CMakeLists.txt
@@ -267,6 +267,7 @@ set(SYCL_COMMON_SOURCES
     "detail/device_filter.cpp"
     "detail/host_pipe_map.cpp"
     "detail/device_global_map.cpp"
+    "detail/kernel_global_info.cpp"
     "detail/device_global_map_entry.cpp"
     "detail/device_image_impl.cpp"
     "detail/device_impl.cpp"
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index de5daeda0c0ab..8fd414ad35da3 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -27,7 +27,6 @@
 #include <cstdint>
 #include <cstring>
 #include <memory>
-#include <unordered_map>
 #include <unordered_set>
 #include <vector>
 
@@ -1048,18 +1047,15 @@ class kernel_bundle_impl
 
   DeviceGlobalMap &getDeviceGlobalMap() { return MDeviceGlobals; }
 
-  void AddKernelArgsSize(const std::string &KernelName, unsigned Size) {
-    auto It = MFreeFuncKernelArgsSizeMap.find(KernelName);
-    if (It == MFreeFuncKernelArgsSizeMap.end()) {
-      MFreeFuncKernelArgsSizeMap[KernelName] = Size;
-    }
-  }
-
   unsigned GetKernelArgsSize(const std::string &KernelName) const {
-    auto It = MFreeFuncKernelArgsSizeMap.find(KernelName);
-    if (It == MFreeFuncKernelArgsSizeMap.end())
+    auto &PM = sycl::detail::ProgramManager::getInstance();
+    const void *GlobalPointer = PM.getKernelGLobalInfoDesc(KernelName.c_str());
+    if (!GlobalPointer)
+      return 0;
+    const unsigned *SizePtr = reinterpret_cast<const unsigned *>(GlobalPointer);
+    if (!SizePtr)
       return 0;
-    return It->second;
+    return *SizePtr;
   }
 
 private:
@@ -1124,7 +1120,6 @@ class kernel_bundle_impl
 
   context MContext;
   std::vector<device_impl *> MDevices;
-  std::unordered_map<std::string, unsigned> MFreeFuncKernelArgsSizeMap;
 
   // For sycl_jit, building from source may have produced sycl binaries that
   // the kernel_bundles now manage.
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
new file mode 100644
index 0000000000000..ad38c5d36d734
--- /dev/null
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -0,0 +1,22 @@
+//==-------------------- kernel_global_info.cpp -----------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <detail/program_manager/program_manager.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail::free_function_info_map {
+
+__SYCL_EXPORT void add(const void *KernelGlobalPtr, const char *UniqueId) {
+  detail::ProgramManager::getInstance().addOrInitKernelGlobalInfo(
+      KernelGlobalPtr, UniqueId);
+}
+
+} // namespace detail::free_function_info_map
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index e5cfa7da08010..4717065adf1e1 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2407,6 +2407,19 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
 }
 
+void ProgramManager::addOrInitKernelGlobalInfo(const void *KernelGlobalPtr,
+                                               const char *UniqueId) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
+  m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
+}
+
+const void *ProgramManager::getKernelGLobalInfoDesc(const char *UniqueId) {
+  const auto It = m_KernelGlobalInfo.find(UniqueId);
+  if (It == m_KernelGlobalInfo.end())
+    return nullptr;
+  return It->second;
+}
+
 std::set<const RTDeviceBinaryImage *>
 ProgramManager::getRawDeviceImages(const std::vector<kernel_id> &KernelIDs) {
   std::set<const RTDeviceBinaryImage *> BinImages;
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index b7b0475457cfa..2534c8115a9a4 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -255,6 +255,14 @@ class ProgramManager {
   void addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
                                   const char *UniqueId);
 
+  // The function inserts or initializes a kernel global desc into the
+  // kernel global map.
+  void addOrInitKernelGlobalInfo(const void *KernelGlobalPtr, const char *UniqueId);
+
+  // The function returns a pointer to the kernel global desc identified by
+  // the unique ID from the kernel global map.
+  const void *getKernelGLobalInfoDesc(const char *UniqueId);
+
   // Returns true if any available image is compatible with the device Dev.
   bool hasCompatibleImage(const device_impl &DeviceImpl);
 
@@ -419,7 +427,7 @@ class ProgramManager {
   bool isBfloat16DeviceImage(const RTDeviceBinaryImage *BinImage);
   bool shouldBF16DeviceImageBeUsed(const RTDeviceBinaryImage *BinImage,
                                    const device_impl &DeviceImpl);
-
+  
 protected:
   /// The three maps below are used during kernel resolution. Any kernel is
   /// identified by its name.
@@ -538,6 +546,10 @@ class ProgramManager {
   // their associated entry resources when they die.
   DeviceGlobalMap m_DeviceGlobals{/*OwnerControlledCleanup=*/true};
 
+  // Maps between host_pipe identifiers and associated kernel global
+  // information.
+  std::unordered_map<KernelNameStrT, const void *> m_KernelGlobalInfo;
+
   // Maps between host_pipe identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>
       m_HostPipes;
diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
index dc1fdd5ea6790..3476c8747102f 100644
--- a/sycl/source/kernel_bundle.cpp
+++ b/sycl/source/kernel_bundle.cpp
@@ -162,11 +162,6 @@ size_t kernel_bundle_plain::ext_oneapi_get_device_global_size(
       std::string(std::string_view(name)));
 }
 
-void kernel_bundle_plain::addFreeFuncKernelArgsSize(unsigned Size,
-                                                    const kernel_id &KernelID) {
-  impl->AddKernelArgsSize(KernelID.get_name(), Size);
-}
-
 //////////////////////////////////
 ///// sycl::detail free functions
 //////////////////////////////////
diff --git a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
index 68e57fe534c8e..4a8a8cc043c5e 100644
--- a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
@@ -32,27 +32,42 @@ int test_num_args_free_function_api(sycl::context &ctxt, sycl::device &dev,
   const bool res = actual == expected_num_args;
   if (!res)
     std::cout << FFTestMark << "test_num_args failed: expected_num_args "
-              << expected_num_args << "actual " << actual << std::endl;
+              << expected_num_args << " actual " << actual << std::endl;
   return res ? 0 : 1;
 }
 
 template <auto *Func>
-int test_num_args_kernel_api(sycl::context &ctxt, sycl::device &dev,
-                             const int expected_num_args) {
+int test_num_args_kernel_api(sycl::context &ctxt, const int expected_num_args) {
   auto bundle =
       syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
   const int actual = bundle.template ext_oneapi_get_kernel<Func>()
                          .template get_info<sycl::info::kernel::num_args>();
-  std::cout << FFTestMark << "actual number of args: " << actual
-            << " expected: " << expected_num_args << std::endl;
   const bool res = actual == expected_num_args;
   if (!res)
     std::cout << FFTestMark
               << "test_num_args_kernel_api failed: expected_num_args "
-              << expected_num_args << "actual " << actual << std::endl;
+              << expected_num_args << " actual " << actual << std::endl;
   return res ? 0 : 1;
 }
 
+template <auto *Func>
+int test_num_args_kernel_id(sycl::context &ctxt, const int expected_num_args) {
+  auto KernelId = syclexp::get_kernel_id<Func>();
+  auto Bundle =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  if (Bundle.has_kernel(KernelId)) {
+    sycl::kernel Kernel = Bundle.get_kernel(KernelId);
+    unsigned actual = Kernel.get_info<sycl::info::kernel::num_args>();
+    const bool res = actual == expected_num_args;
+    if (!res)
+      std::cout << FFTestMark
+                << "test_num_args_kernel_id failed: expected_num_args "
+                << expected_num_args << " actual " << actual << std::endl;
+    return res ? 0 : 1;
+  }
+  return 1;
+}
+
 int main() {
   sycl::queue q;
   sycl::context ctx = q.get_context();
@@ -61,8 +76,12 @@ int main() {
   int ret = test_num_args_free_function_api<func_range>(ctx, dev, 2);
   ret |= test_num_args_free_function_api<func_single>(ctx, dev, 2);
   ret |= test_num_args_free_function_api<kernel_func>(ctx, dev, 3);
-  ret |= test_num_args_kernel_api<func_range>(ctx, dev, 2);
-  ret |= test_num_args_kernel_api<func_single>(ctx, dev, 2);
-  ret |= test_num_args_kernel_api<kernel_func>(ctx, dev, 3);
+  ret |= test_num_args_kernel_api<func_range>(ctx, 2);
+  ret |= test_num_args_kernel_api<func_single>(ctx, 2);
+  ret |= test_num_args_kernel_api<kernel_func>(ctx, 3);
+  ret |= test_num_args_kernel_id<func_range>(ctx, 2);
+  ret |= test_num_args_kernel_id<func_single>(ctx, 2);
+  ret |= test_num_args_kernel_id<kernel_func>(ctx, 3);
+
   return ret;
 }
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 4708e32ca4df1..62807bae007a7 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3842,10 +3842,10 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain33contains_specialization_constantsEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain3endEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5beginEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
-_ZN4sycl3_V16detail19kernel_bundle_plain25addFreeFuncKernelArgsSizeEjRKNS0_9kernel_idE
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
+_ZN4sycl3_V16detail22free_function_info_map3addEPKvPKc
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost10getSamplerEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost12getNumOfDimsEv

From bf4786efd8e686bc73197c50f8fb0bb9b09196a5 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 14:52:45 +0200
Subject: [PATCH 12/23] [SYCL] do not create a separate function  to add global
 kernel descs in header [SYCL] fix formatting

---
 clang/lib/Sema/SemaSYCL.cpp                   |  30 +--
 ...ee_function_default_template_arguments.cpp |  15 ++
 .../CodeGenSYCL/free_function_int_header.cpp  | 254 +-----------------
 .../free_function_int_header_rtc_mode.cpp     |  20 +-
 .../sycl/detail/kernel_global_info.hpp        |   4 +-
 sycl/include/sycl/kernel_bundle.hpp           |   8 +
 sycl/source/detail/kernel_bundle_impl.hpp     |   2 +-
 sycl/source/detail/kernel_global_info.cpp     |   4 +-
 .../program_manager/program_manager.cpp       |   7 +-
 .../program_manager/program_manager.hpp       |   7 +-
 .../test-e2e/FreeFunctionKernels/num_args.cpp |   5 +-
 11 files changed, 68 insertions(+), 288 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 91d7e54570889..0b1b2c8dd3273 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7191,31 +7191,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   if (FreeFunctionCount > 0) {
     O << "\n#include <sycl/kernel_bundle.hpp>\n";
   }
-  ShimCounter = 1;
-  for (const KernelDesc &K : KernelDescs) {
-    if (!S.isFreeFunction(K.SyclKernel))
-      continue;
-
-    O << "\n// Definition of kernel_id of " << K.Name << "\n";
-    O << "namespace sycl {\n";
-    O << "template <>\n";
-    O << "inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim"
-      << ShimCounter << "()>() {\n";
-    O << "  return sycl::detail::get_kernel_id_impl(std::string_view{"
-      << "sycl::detail::FreeFunctionInfoData<__sycl_shim" << ShimCounter
-      << "()>::getFunctionName()});\n";
-    O << "}\n";
-    O << "}\n";
-    ++ShimCounter;
-  }
   O << "#include <sycl/detail/kernel_global_info.hpp>\n";
 
   ShimCounter = 0;
   O << "namespace sycl {\n";
   O << "inline namespace _V1 {\n";
   O << "namespace detail {\n";
-  O << "namespace free_function_map {\n";
-  O << "inline void update_device_global_map() {\n";
+  O << "struct GlobalMapUpdater {\n";
+  O << "  GlobalMapUpdater() {\n";
   for (const KernelDesc &K : KernelDescs) {
     if (!S.isFreeFunction(K.SyclKernel))
       continue;
@@ -7225,18 +7208,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       << "]);\n";
     ++ShimCounter;
   }
-  O << "}\n";
-  O << "struct GlobalMapUpdater {\n";
-  O << "  GlobalMapUpdater() {\n";
-  O << "    update_device_global_map();\n";
-  O << "    std::cout << \"Device global map updated for free function "
-       "kernels. "
-       "Total: "
-    << ShimCounter << " kernels.\" << std::endl;\n";
   O << "  }\n";
   O << "};\n";
   O << "static GlobalMapUpdater updater;\n";
-  O << "} // namespace free_function_map\n";
   O << "} // namespace detail\n";
   O << "} // namespace _V1\n";
   O << "} // namespace sycl\n";
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 45ba5c8026f07..515255d9da8e4 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -1084,3 +1084,18 @@ namespace Testing::Tests {
 // CHECK-NEXT:    static constexpr bool value = true;
 // CHECK-NEXT:  };
 // CHECK-NEXT:  }
+
+// CHECK: #include <sycl/kernel_bundle.hpp>
+// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: struct GlobalMapUpdater {
+// CHECK-NEXT:   GlobalMapUpdater() {
+// CHECK-COUNT-28: sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes +
+// CHECK-NEXT:   }
+// CHECK-NEXT: };
+// CHECK-NEXT: static GlobalMapUpdater updater;
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index 34ebc6668558b..d17881abffc7e 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1564,248 +1564,18 @@ void ff_24(int arg) {
 
 // CHECK-NEXT: }
 
-// CHECK: #include <sycl/kernel_bundle.hpp>
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:   template <>
-// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim1()>() {
-// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim1()>::getFunctionName()});
-// CHECK-NEXT:   }
-// CHECK-NEXT: }
 
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piiii
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:   template <>
-// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim2()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});
-// CHECK-NEXT:   }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IiEvPT_S0_S0_
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:   template <>
-// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim3()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim3()>::getFunctionName()});
-// CHECK-NEXT:   }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IfEvPT_S0_S0_
+// CHECK: #include <sycl/kernel_bundle.hpp>
+// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
 // CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:   template <>
-// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim4()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim4()>::getFunctionName()});
+// CHECK-NEXT: inline namespace _V1 {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: struct GlobalMapUpdater {
+// CHECK-NEXT:  GlobalMapUpdater() {
+// CHECK-COUNT-33: sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 
 // CHECK-NEXT:   }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IdEvPT_S0_S0_
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:   template <>
-// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim5()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim5()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim6()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim6()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim7()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim7()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim8()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim9()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim9()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-
-// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions4ff_9EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim10()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim11()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim11()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  } 
-
-// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim12()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim12()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: Definition of kernel_id of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim13()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim13()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_13EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim14()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim14()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim15()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim15()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-//
-// CHECK: // Definition of kernel_id of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim16()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim16()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_11N4sycl3_V114local_accessorIiLi1EEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim17()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim17()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_11IfEvN4sycl3_V114local_accessorIT_Li1EEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim18()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim18()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_12N4sycl3_V17samplerE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim19()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim19()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_13N4sycl3_V16streamE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim20()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim20()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_14N4sycl3_V13ext6oneapi12experimental13annotated_argIiJEEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim21()>() {
-// CHECK-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim21()>::getFunctionName()});
-
-// CHECK: Definition of kernel_id of _Z19__sycl_kernel_ff_15N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim22()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim22()>::getFunctionName()});
-
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim23()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim23()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-  
-// CHECK:  // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_15EiPi
-// CHECK-NEXT:  namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim24()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim24()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-  
-// CHECK:  // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_
-// CHECK-NEXT:  namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim25()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim25()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim26()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim26()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT:  template <>
-// CHECK-NEXT:  kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim27()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim27()>::getFunctionName()});
-// CHECK-NEXT:  }
-// CHECK-NEXT:  }
-
-// CHECK: // Definition of kernel_id of _Z19__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim29()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim29()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim30()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim30()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim31()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim31()>::getFunctionName()});
-
-// CHECK: // Definition of kernel_id of _Z19__sycl_kernel_ff_24i
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim32()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim32()>::getFunctionName()});
-// CHECK-NEXT: }
-// CHECK-NEXT: }
-
-// CHECK: // Definition of kernel_id of _Z19__sycl_kernel_ff_23i
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: template <>
-// CHECK-NEXT: inline kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim33()>() {
-// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim33()>::getFunctionName()});
-
-// CHECK-NEXT: }
-// CHECK-NEXT: }
+// CHECK-NEXT: };
+// CHECK-NEXT: static GlobalMapUpdater updater;
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace _V1
+// CHECK-NEXT: } // namespace sycl
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index fd0016846e568..0b91b364336f7 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -73,7 +73,19 @@ int main(){
 // CHECK-NORTC:       struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim2(), 2> {
 // CHECK-NORTC-NEXT:  static constexpr bool value = true;
 
-// CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() {
-// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim1()>::getFunctionName()});
-// CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() {
-// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{sycl::detail::FreeFunctionInfoData<__sycl_shim2()>::getFunctionName()});
+
+// CHECK-NORTC: #include <sycl/kernel_bundle.hpp>
+// CHECK-NORTC-NEXT: #include <sycl/detail/kernel_global_info.hpp>
+// CHECK-NORTC-NEXT: namespace sycl {
+// CHECK-NORTC-NEXT: inline namespace _V1 {
+// CHECK-NORTC-NEXT: namespace detail {
+// CHECK-NORTC-NEXT: struct GlobalMapUpdater {
+// CHECK-NORTC-NEXT:   GlobalMapUpdater() {
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 0), sycl::detail::kernel_names[0]);
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 1), sycl::detail::kernel_names[1]);
+// CHECK-NORTC-NEXT:   }
+// CHECK-NORTC-NEXT: };
+// CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
+// CHECK-NORTC-NEXT: } // namespace detail
+// CHECK-NORTC-NEXT: } // namespace _V1
+// CHECK-NORTC-NEXT: } // namespace sycl
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
index a6b6f1593f0b9..9903e088c81cc 100644
--- a/sycl/include/sycl/detail/kernel_global_info.hpp
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -1,4 +1,4 @@
-//==-------------------- kernel_global_info.hpp -----------------------------==//
+//==-------------------- kernel_global_info.hpp -----------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -20,4 +20,4 @@ __SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
 } // namespace free_function_info_map
 } // namespace detail
 } // namespace _V1
-} // namespace sycl
\ No newline at end of file
+} // namespace sycl
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index fbe7024d412a1..bc57a64c6e0dd 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -291,6 +291,7 @@ class __SYCL_EXPORT kernel_bundle_plain {
   // \returns a bool value which indicates if specialization constant was set to
   // a value different from default value.
   bool is_specialization_constant_set(const char *SpecName) const noexcept;
+
   detail::KernelBundleImplPtr impl;
 
 private:
@@ -729,6 +730,13 @@ get_kernel_bundle(const context &Ctx) {
   return get_kernel_bundle<State>(Ctx, Ctx.get_devices(),
                                   {get_kernel_id<Func>()});
 }
+
+template <auto *Func>
+std::enable_if_t<is_kernel_v<Func>, kernel_id> get_kernel_id() {
+  return get_kernel_id_impl(detail::string_view(
+      detail::FreeFunctionInfoData<Func>::getFunctionName()));
+}
+
 } // namespace ext::oneapi::experimental
 
 namespace detail {
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 8fd414ad35da3..3e3373341a9b2 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -1049,7 +1049,7 @@ class kernel_bundle_impl
 
   unsigned GetKernelArgsSize(const std::string &KernelName) const {
     auto &PM = sycl::detail::ProgramManager::getInstance();
-    const void *GlobalPointer = PM.getKernelGLobalInfoDesc(KernelName.c_str());
+    const void *GlobalPointer = PM.getKernelGlobalInfoDesc(KernelName.c_str());
     if (!GlobalPointer)
       return 0;
     const unsigned *SizePtr = reinterpret_cast<const unsigned *>(GlobalPointer);
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index ad38c5d36d734..6743ab76132ce 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -1,4 +1,4 @@
-//==-------------------- kernel_global_info.cpp -----------------------------==//
+//==-------------------- kernel_global_info.cpp ----------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -13,7 +13,7 @@ inline namespace _V1 {
 namespace detail::free_function_info_map {
 
 __SYCL_EXPORT void add(const void *KernelGlobalPtr, const char *UniqueId) {
-  detail::ProgramManager::getInstance().addOrInitKernelGlobalInfo(
+  detail::ProgramManager::getInstance().registerKernelGlobalInfo(
       KernelGlobalPtr, UniqueId);
 }
 
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 4717065adf1e1..3741207633503 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2407,13 +2407,14 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
 }
 
-void ProgramManager::addOrInitKernelGlobalInfo(const void *KernelGlobalPtr,
-                                               const char *UniqueId) {
+void ProgramManager::registerKernelGlobalInfo(const void *KernelGlobalPtr,
+                                              const char *UniqueId) {
   std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
 }
 
-const void *ProgramManager::getKernelGLobalInfoDesc(const char *UniqueId) {
+const void *ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   const auto It = m_KernelGlobalInfo.find(UniqueId);
   if (It == m_KernelGlobalInfo.end())
     return nullptr;
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 2534c8115a9a4..c4602bc022440 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -257,11 +257,12 @@ class ProgramManager {
 
   // The function inserts or initializes a kernel global desc into the
   // kernel global map.
-  void addOrInitKernelGlobalInfo(const void *KernelGlobalPtr, const char *UniqueId);
+  void registerKernelGlobalInfo(const void *KernelGlobalPtr,
+                                const char *UniqueId);
 
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
-  const void *getKernelGLobalInfoDesc(const char *UniqueId);
+  const void *getKernelGlobalInfoDesc(const char *UniqueId);
 
   // Returns true if any available image is compatible with the device Dev.
   bool hasCompatibleImage(const device_impl &DeviceImpl);
@@ -427,7 +428,7 @@ class ProgramManager {
   bool isBfloat16DeviceImage(const RTDeviceBinaryImage *BinImage);
   bool shouldBF16DeviceImageBeUsed(const RTDeviceBinaryImage *BinImage,
                                    const device_impl &DeviceImpl);
-  
+
 protected:
   /// The three maps below are used during kernel resolution. Any kernel is
   /// identified by its name.
diff --git a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
index 4a8a8cc043c5e..701b900036a62 100644
--- a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
@@ -1,6 +1,5 @@
-// REQUIRES: level_zero, level_zero_dev_kit
-// RUN: %{build} %level_zero_options -o %t.ze.out
-// RUN: %{run} %t.ze.out
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
 
 #include <iostream>
 #include <sycl/detail/core.hpp>

From 811c4b6d898ef72b28d2bb38560cbd78c6946539 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 15:58:20 +0200
Subject: [PATCH 13/23] [SYCL] do not use string pattern to detect if kernel is
 a free function

---
 sycl/source/detail/kernel_bundle_impl.hpp |  7 ++++---
 sycl/source/detail/kernel_impl.cpp        | 14 ++------------
 sycl/source/detail/kernel_impl.hpp        |  7 +++----
 3 files changed, 9 insertions(+), 19 deletions(-)

diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 3e3373341a9b2..b07ca8ad214c6 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -1047,14 +1047,15 @@ class kernel_bundle_impl
 
   DeviceGlobalMap &getDeviceGlobalMap() { return MDeviceGlobals; }
 
-  unsigned GetKernelArgsSize(const std::string &KernelName) const {
+  std::optional<unsigned>
+  tryGetKernelArgsSize(const std::string &KernelName) const {
     auto &PM = sycl::detail::ProgramManager::getInstance();
     const void *GlobalPointer = PM.getKernelGlobalInfoDesc(KernelName.c_str());
     if (!GlobalPointer)
-      return 0;
+      return std::nullopt;
     const unsigned *SizePtr = reinterpret_cast<const unsigned *>(GlobalPointer);
     if (!SizePtr)
-      return 0;
+      return std::nullopt;
     return *SizePtr;
   }
 
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index bccac77366679..4926d59d89e0c 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -37,9 +37,6 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &Context,
 
   // Enable USM indirect access for interoperability kernels.
   enableUSMIndirectAccess();
-  const std::string KernelName = get_info<info::kernel::function_name>();
-  const auto pos = KernelName.find("__sycl_kernel_");
-  isFreeFuncKernel = pos != std::string::npos;
 }
 
 kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
@@ -59,10 +56,6 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
   // path.
   if (MCreatedFromSource || MIsInterop)
     enableUSMIndirectAccess();
-
-  const std::string KernelName = get_info<info::kernel::function_name>();
-  const auto pos = KernelName.find("__sycl_kernel_");
-  isFreeFuncKernel = pos != std::string::npos;
 }
 
 kernel_impl::~kernel_impl() {
@@ -129,9 +122,6 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       }))
     return;
 
-  if (isFreeFuncKernel)
-    return;
-
   throw sycl::exception(
       sycl::make_error_code(errc::invalid),
       "info::kernel::num_args descriptor may only be used to query a kernel "
@@ -139,9 +129,9 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       "interoperability function or to query a device built-in kernel");
 }
 
-unsigned kernel_impl ::getFreeFuncKernelArgSize() const {
+std::optional<unsigned> kernel_impl ::getFreeFuncKernelArgSize() const {
   const std::string KernelName = get_info<info::kernel::function_name>();
-  return MKernelBundleImpl->GetKernelArgsSize(KernelName);
+  return MKernelBundleImpl->tryGetKernelArgsSize(KernelName);
 }
 
 void kernel_impl::enableUSMIndirectAccess() const {
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index 6ee9e74d57c28..32c52309f1d80 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -254,7 +254,6 @@ class kernel_impl {
   const KernelArgMask *MKernelArgMaskPtr;
   std::mutex *MCacheMutex = nullptr;
   mutable std::string MName;
-  bool isFreeFuncKernel = false;
 
   bool isBuiltInKernel(device_impl &Device) const;
   void checkIfValidForNumArgsInfoQuery() const;
@@ -271,7 +270,7 @@ class kernel_impl {
                                size_t DynamicLocalMemorySize) const;
 
   void enableUSMIndirectAccess() const;
-  unsigned getFreeFuncKernelArgSize() const;
+  std::optional<unsigned> getFreeFuncKernelArgSize() const;
 };
 
 template <int Dimensions>
@@ -312,9 +311,9 @@ inline typename Param::return_type kernel_impl::get_info() const {
   static_assert(is_kernel_info_desc<Param>::value,
                 "Invalid kernel information descriptor");
   if constexpr (std::is_same_v<Param, info::kernel::num_args>) {
+    if (std::optional<unsigned> FFArgSize = getFreeFuncKernelArgSize())
+      return *FFArgSize;
     checkIfValidForNumArgsInfoQuery();
-    if (isFreeFuncKernel)
-      return getFreeFuncKernelArgSize();
   }
   return get_kernel_info<Param>(this->getHandleRef(), getAdapter());
 }

From 618f51eb655237de4c92b2b068ff62cfbf241ddd Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 15:58:44 +0200
Subject: [PATCH 14/23] [SYCL][TEST] update win abi test

---
 sycl/test/abi/sycl_symbols_windows.dump | 1 +
 1 file changed, 1 insertion(+)

diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index e87045fc6ae1f..62a1e1b3ecc80 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -3765,6 +3765,7 @@
 ?__trunc_impl@_V1@sycl@@YANN@Z
 ?accelerator_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z
 ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
+?add@free_function_info_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA?AVnode@34567@AEBVproperty_list@67@@Z
 ?addAccessorReq@handler@_V1@sycl@@AEAAXV?$shared_ptr@VAccessorImplHost@detail@_V1@sycl@@@std@@@Z

From ef30cdedad0aaa16e5e4a05ef6cd3505a37141f2 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 17:35:14 +0200
Subject: [PATCH 15/23] [SYCL][TEST] add new headers to test input directory

---
 .../SemaSYCL/Inputs/sycl/detail/export.hpp    | 39 +++++++++++++++++++
 .../Inputs/sycl/detail/kernel_global_info.hpp | 23 +++++++++++
 2 files changed, 62 insertions(+)
 create mode 100644 clang/test/SemaSYCL/Inputs/sycl/detail/export.hpp
 create mode 100644 clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp

diff --git a/clang/test/SemaSYCL/Inputs/sycl/detail/export.hpp b/clang/test/SemaSYCL/Inputs/sycl/detail/export.hpp
new file mode 100644
index 0000000000000..98b5673a079e0
--- /dev/null
+++ b/clang/test/SemaSYCL/Inputs/sycl/detail/export.hpp
@@ -0,0 +1,39 @@
+//==---------------- export.hpp - SYCL standard header file ----------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifndef __SYCL_DEVICE_ONLY__
+#ifndef __SYCL_EXPORT
+#ifdef _WIN32
+
+#define __SYCL_DLL_LOCAL
+
+#if __SYCL_BUILD_SYCL_DLL
+#define __SYCL_EXPORT __declspec(dllexport)
+#define __SYCL_EXPORT_DEPRECATED(x) __declspec(dllexport, deprecated(x))
+#else
+#define __SYCL_EXPORT __declspec(dllimport)
+#define __SYCL_EXPORT_DEPRECATED(x) __declspec(dllimport, deprecated(x))
+#endif //__SYCL_BUILD_SYCL_DLL
+#else  // _WIN32
+
+#define __SYCL_DLL_LOCAL __attribute__((visibility("hidden")))
+
+#define __SYCL_EXPORT __attribute__((visibility("default")))
+#define __SYCL_EXPORT_DEPRECATED(x)                                            \
+  __attribute__((visibility("default"), deprecated(x)))
+#endif // _WIN32
+#endif // __SYCL_EXPORT
+#else
+#ifndef __SYCL_EXPORT
+#define __SYCL_EXPORT
+#define __SYCL_EXPORT_DEPRECATED(x)
+#define __SYCL_DLL_LOCAL
+#endif
+#endif // __SYCL_DEVICE_ONLY__
diff --git a/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp b/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
new file mode 100644
index 0000000000000..9903e088c81cc
--- /dev/null
+++ b/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
@@ -0,0 +1,23 @@
+//==-------------------- kernel_global_info.hpp -----------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/export.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+namespace free_function_info_map {
+
+__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
+
+} // namespace free_function_info_map
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

From ed5acdf11f6f108534e142b15472b26aa3dcfcc8 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 17:50:15 +0200
Subject: [PATCH 16/23] [SYCL] update args to match map

---
 clang/lib/Sema/SemaSYCL.cpp                                 | 6 +++---
 .../free_function_default_template_arguments.cpp            | 2 +-
 clang/test/CodeGenSYCL/free_function_int_header.cpp         | 2 +-
 .../test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp  | 4 ++--
 sycl/include/sycl/detail/kernel_global_info.hpp             | 2 +-
 sycl/source/detail/kernel_global_info.cpp                   | 4 ++--
 sycl/source/detail/program_manager/program_manager.cpp      | 4 ++--
 sycl/source/detail/program_manager/program_manager.hpp      | 4 ++--
 sycl/test/abi/sycl_symbols_linux.dump                       | 2 +-
 9 files changed, 15 insertions(+), 15 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 0b1b2c8dd3273..efa7acbe2dc45 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7203,9 +7203,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     if (!S.isFreeFunction(K.SyclKernel))
       continue;
     O << "sycl::detail::free_function_info_map::add("
-      << "reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + "
-      << ShimCounter << "), sycl::detail::kernel_names[" << ShimCounter
-      << "]);\n";
+      << "sycl::detail::kernel_names[" << ShimCounter
+      << "], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + "
+      << ShimCounter << "));\n";
     ++ShimCounter;
   }
   O << "  }\n";
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 515255d9da8e4..39f615a6bc248 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -1092,7 +1092,7 @@ namespace Testing::Tests {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:   GlobalMapUpdater() {
-// CHECK-COUNT-28: sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes +
+// CHECK-COUNT-28: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index d17881abffc7e..0325dd5645ea3 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1572,7 +1572,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:  GlobalMapUpdater() {
-// CHECK-COUNT-33: sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 
+// CHECK-COUNT-33: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[ 
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 0b91b364336f7..edd14f84e8cfd 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -81,8 +81,8 @@ int main(){
 // CHECK-NORTC-NEXT: namespace detail {
 // CHECK-NORTC-NEXT: struct GlobalMapUpdater {
 // CHECK-NORTC-NEXT:   GlobalMapUpdater() {
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 0), sycl::detail::kernel_names[0]);
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 1), sycl::detail::kernel_names[1]);
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[0], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 0));
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[1], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 1));
 // CHECK-NORTC-NEXT:   }
 // CHECK-NORTC-NEXT: };
 // CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
index 9903e088c81cc..078a4a50e8eed 100644
--- a/sycl/include/sycl/detail/kernel_global_info.hpp
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -15,7 +15,7 @@ inline namespace _V1 {
 namespace detail {
 namespace free_function_info_map {
 
-__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
+__SYCL_EXPORT void add(const char *UniqueId, const void *DeviceGlobalPtr);
 
 } // namespace free_function_info_map
 } // namespace detail
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index 6743ab76132ce..49fb5b92dde8f 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -12,9 +12,9 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail::free_function_info_map {
 
-__SYCL_EXPORT void add(const void *KernelGlobalPtr, const char *UniqueId) {
+__SYCL_EXPORT void add(const char *UniqueId, const void *KernelGlobalPtr) {
   detail::ProgramManager::getInstance().registerKernelGlobalInfo(
-      KernelGlobalPtr, UniqueId);
+      UniqueId, KernelGlobalPtr);
 }
 
 } // namespace detail::free_function_info_map
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 3741207633503..5e54129e032ff 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2407,8 +2407,8 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
 }
 
-void ProgramManager::registerKernelGlobalInfo(const void *KernelGlobalPtr,
-                                              const char *UniqueId) {
+void ProgramManager::registerKernelGlobalInfo(const char *UniqueId,
+                                              const void *KernelGlobalPtr) {
   std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
 }
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index c4602bc022440..6e4e586874fb1 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -257,8 +257,8 @@ class ProgramManager {
 
   // The function inserts or initializes a kernel global desc into the
   // kernel global map.
-  void registerKernelGlobalInfo(const void *KernelGlobalPtr,
-                                const char *UniqueId);
+  void registerKernelGlobalInfo(const char *UniqueId,
+                                const void *KernelGlobalPtr);
 
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 62807bae007a7..2b2779279239f 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3845,7 +3845,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
-_ZN4sycl3_V16detail22free_function_info_map3addEPKvPKc
+_ZN4sycl3_V16detail22free_function_info_map3addEPKcPKv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost10getSamplerEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost12getNumOfDimsEv

From b05235948d8e4d9897a732b657a5925b5070101f Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 18:27:57 +0200
Subject: [PATCH 17/23] [SYCL] do not lock if add or get kernel global info

---
 sycl/source/detail/kernel_bundle_impl.hpp              | 9 +++------
 sycl/source/detail/kernel_impl.cpp                     | 3 +--
 sycl/source/detail/program_manager/program_manager.cpp | 2 --
 3 files changed, 4 insertions(+), 10 deletions(-)

diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index b07ca8ad214c6..4598120fd723f 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -1048,15 +1048,12 @@ class kernel_bundle_impl
   DeviceGlobalMap &getDeviceGlobalMap() { return MDeviceGlobals; }
 
   std::optional<unsigned>
-  tryGetKernelArgsSize(const std::string &KernelName) const {
+  tryGetKernelArgsSize(const std::string_view KernelName) const {
     auto &PM = sycl::detail::ProgramManager::getInstance();
-    const void *GlobalPointer = PM.getKernelGlobalInfoDesc(KernelName.c_str());
+    const void *GlobalPointer = PM.getKernelGlobalInfoDesc(KernelName.data());
     if (!GlobalPointer)
       return std::nullopt;
-    const unsigned *SizePtr = reinterpret_cast<const unsigned *>(GlobalPointer);
-    if (!SizePtr)
-      return std::nullopt;
-    return *SizePtr;
+    return *reinterpret_cast<const unsigned *>(GlobalPointer);
   }
 
 private:
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 4926d59d89e0c..d82e93579f4aa 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -130,8 +130,7 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
 }
 
 std::optional<unsigned> kernel_impl ::getFreeFuncKernelArgSize() const {
-  const std::string KernelName = get_info<info::kernel::function_name>();
-  return MKernelBundleImpl->tryGetKernelArgsSize(KernelName);
+  return MKernelBundleImpl->tryGetKernelArgsSize(getName());
 }
 
 void kernel_impl::enableUSMIndirectAccess() const {
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 5e54129e032ff..61cc95f43e65c 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2409,12 +2409,10 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
 
 void ProgramManager::registerKernelGlobalInfo(const char *UniqueId,
                                               const void *KernelGlobalPtr) {
-  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
 }
 
 const void *ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
-  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   const auto It = m_KernelGlobalInfo.find(UniqueId);
   if (It == m_KernelGlobalInfo.end())
     return nullptr;

From 0a35b308963641dd35b6eb25746b6137f3e8ef93 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Wed, 23 Jul 2025 19:26:50 +0200
Subject: [PATCH 18/23] [SYCL][TEST] update win abi test

---
 sycl/test/abi/sycl_symbols_windows.dump | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index 62a1e1b3ecc80..d359b9e5c8dd2 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -3765,7 +3765,7 @@
 ?__trunc_impl@_V1@sycl@@YANN@Z
 ?accelerator_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z
 ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
-?add@free_function_info_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
+?add@free_function_info_map@detail@_V1@sycl@@YAXPEBDPEBX@Z
 ?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA?AVnode@34567@AEBVproperty_list@67@@Z
 ?addAccessorReq@handler@_V1@sycl@@AEAAXV?$shared_ptr@VAccessorImplHost@detail@_V1@sycl@@@std@@@Z

From f4efc2f18de432b7663bf99f9f48ddf40086f216 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Thu, 24 Jul 2025 10:07:14 +0200
Subject: [PATCH 19/23] [SYCL] do not cast unsigned to void pointer

---
 clang/lib/Sema/SemaSYCL.cpp                                | 3 +--
 .../test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp | 4 ++--
 sycl/include/sycl/detail/kernel_global_info.hpp            | 2 +-
 sycl/source/detail/kernel_bundle_impl.hpp                  | 5 +----
 sycl/source/detail/kernel_global_info.cpp                  | 2 +-
 sycl/source/detail/program_manager/program_manager.cpp     | 7 ++++---
 sycl/source/detail/program_manager/program_manager.hpp     | 7 +++----
 sycl/test/abi/sycl_symbols_linux.dump                      | 2 +-
 8 files changed, 14 insertions(+), 18 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index efa7acbe2dc45..ef7b0f1695d19 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7204,8 +7204,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       continue;
     O << "sycl::detail::free_function_info_map::add("
       << "sycl::detail::kernel_names[" << ShimCounter
-      << "], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + "
-      << ShimCounter << "));\n";
+      << "], sycl::detail::kernel_args_sizes[" << ShimCounter << "]);\n";
     ++ShimCounter;
   }
   O << "  }\n";
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index edd14f84e8cfd..34ada95ca19fa 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -81,8 +81,8 @@ int main(){
 // CHECK-NORTC-NEXT: namespace detail {
 // CHECK-NORTC-NEXT: struct GlobalMapUpdater {
 // CHECK-NORTC-NEXT:   GlobalMapUpdater() {
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[0], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 0));
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[1], reinterpret_cast<const void*>(sycl::detail::kernel_args_sizes + 1));
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[0], sycl::detail::kernel_args_sizes[0]);
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[1], sycl::detail::kernel_args_sizes[1]);
 // CHECK-NORTC-NEXT:   }
 // CHECK-NORTC-NEXT: };
 // CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
index 078a4a50e8eed..bef3bb9c35950 100644
--- a/sycl/include/sycl/detail/kernel_global_info.hpp
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -15,7 +15,7 @@ inline namespace _V1 {
 namespace detail {
 namespace free_function_info_map {
 
-__SYCL_EXPORT void add(const char *UniqueId, const void *DeviceGlobalPtr);
+__SYCL_EXPORT void add(const char *UniqueId, unsigned DeviceGlobalPtr);
 
 } // namespace free_function_info_map
 } // namespace detail
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 4598120fd723f..369a535e9552b 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -1050,10 +1050,7 @@ class kernel_bundle_impl
   std::optional<unsigned>
   tryGetKernelArgsSize(const std::string_view KernelName) const {
     auto &PM = sycl::detail::ProgramManager::getInstance();
-    const void *GlobalPointer = PM.getKernelGlobalInfoDesc(KernelName.data());
-    if (!GlobalPointer)
-      return std::nullopt;
-    return *reinterpret_cast<const unsigned *>(GlobalPointer);
+    return PM.getKernelGlobalInfoDesc(KernelName.data());
   }
 
 private:
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index 49fb5b92dde8f..860ec9ece1fed 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -12,7 +12,7 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail::free_function_info_map {
 
-__SYCL_EXPORT void add(const char *UniqueId, const void *KernelGlobalPtr) {
+__SYCL_EXPORT void add(const char *UniqueId, unsigned KernelGlobalPtr) {
   detail::ProgramManager::getInstance().registerKernelGlobalInfo(
       UniqueId, KernelGlobalPtr);
 }
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 61cc95f43e65c..ef4a634766ab5 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2408,14 +2408,15 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
 }
 
 void ProgramManager::registerKernelGlobalInfo(const char *UniqueId,
-                                              const void *KernelGlobalPtr) {
+                                              unsigned KernelGlobalPtr) {
   m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
 }
 
-const void *ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
+std::optional<unsigned>
+ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
   const auto It = m_KernelGlobalInfo.find(UniqueId);
   if (It == m_KernelGlobalInfo.end())
-    return nullptr;
+    return std::nullopt;
   return It->second;
 }
 
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 6e4e586874fb1..029ed38bf88d9 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -257,12 +257,11 @@ class ProgramManager {
 
   // The function inserts or initializes a kernel global desc into the
   // kernel global map.
-  void registerKernelGlobalInfo(const char *UniqueId,
-                                const void *KernelGlobalPtr);
+  void registerKernelGlobalInfo(const char *UniqueId, unsigned KernelGlobalPtr);
 
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
-  const void *getKernelGlobalInfoDesc(const char *UniqueId);
+  std::optional<unsigned> getKernelGlobalInfoDesc(const char *UniqueId);
 
   // Returns true if any available image is compatible with the device Dev.
   bool hasCompatibleImage(const device_impl &DeviceImpl);
@@ -549,7 +548,7 @@ class ProgramManager {
 
   // Maps between host_pipe identifiers and associated kernel global
   // information.
-  std::unordered_map<KernelNameStrT, const void *> m_KernelGlobalInfo;
+  std::unordered_map<KernelNameStrT, unsigned> m_KernelGlobalInfo;
 
   // Maps between host_pipe identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 2b2779279239f..43424e3eb1d82 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3845,7 +3845,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
-_ZN4sycl3_V16detail22free_function_info_map3addEPKcPKv
+_ZN4sycl3_V16detail22free_function_info_map3addEPKcj
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost10getSamplerEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost12getNumOfDimsEv

From 11c62706f83986b601222bf25bfab3b8a843302f Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Thu, 24 Jul 2025 12:51:18 +0200
Subject: [PATCH 20/23] [SYCL] rework register global kernel info function

---
 clang/lib/Sema/SemaSYCL.cpp                   | 35 ++++++++-----------
 ...ee_function_default_template_arguments.cpp |  2 +-
 .../CodeGenSYCL/free_function_int_header.cpp  |  2 +-
 .../free_function_int_header_rtc_mode.cpp     |  3 +-
 .../sycl/detail/kernel_global_info.hpp        |  3 +-
 sycl/source/detail/kernel_global_info.cpp     |  9 +++--
 .../program_manager/program_manager.cpp       | 14 ++++++--
 .../program_manager/program_manager.hpp       |  5 +--
 sycl/test/abi/sycl_symbols_linux.dump         |  2 +-
 9 files changed, 41 insertions(+), 34 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ef7b0f1695d19..857fdf85dd289 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7190,29 +7190,22 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
 
   if (FreeFunctionCount > 0) {
     O << "\n#include <sycl/kernel_bundle.hpp>\n";
-  }
-  O << "#include <sycl/detail/kernel_global_info.hpp>\n";
-
-  ShimCounter = 0;
-  O << "namespace sycl {\n";
-  O << "inline namespace _V1 {\n";
-  O << "namespace detail {\n";
-  O << "struct GlobalMapUpdater {\n";
-  O << "  GlobalMapUpdater() {\n";
-  for (const KernelDesc &K : KernelDescs) {
-    if (!S.isFreeFunction(K.SyclKernel))
-      continue;
+    O << "#include <sycl/detail/kernel_global_info.hpp>\n";
+    O << "namespace sycl {\n";
+    O << "inline namespace _V1 {\n";
+    O << "namespace detail {\n";
+    O << "struct GlobalMapUpdater {\n";
+    O << "  GlobalMapUpdater() {\n";
     O << "sycl::detail::free_function_info_map::add("
-      << "sycl::detail::kernel_names[" << ShimCounter
-      << "], sycl::detail::kernel_args_sizes[" << ShimCounter << "]);\n";
-    ++ShimCounter;
+      << "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
+      << KernelDescs.size() << ");\n";
+    O << "  }\n";
+    O << "};\n";
+    O << "static GlobalMapUpdater updater;\n";
+    O << "} // namespace detail\n";
+    O << "} // namespace _V1\n";
+    O << "} // namespace sycl\n";
   }
-  O << "  }\n";
-  O << "};\n";
-  O << "static GlobalMapUpdater updater;\n";
-  O << "} // namespace detail\n";
-  O << "} // namespace _V1\n";
-  O << "} // namespace sycl\n";
 }
 
 bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) {
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 39f615a6bc248..dc542cb50f389 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -1092,7 +1092,7 @@ namespace Testing::Tests {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:   GlobalMapUpdater() {
-// CHECK-COUNT-28: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[
+// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index 0325dd5645ea3..add97822677f4 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1572,7 +1572,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:  GlobalMapUpdater() {
-// CHECK-COUNT-33: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[ 
+// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 34ada95ca19fa..0009e71c7c9f1 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -81,8 +81,7 @@ int main(){
 // CHECK-NORTC-NEXT: namespace detail {
 // CHECK-NORTC-NEXT: struct GlobalMapUpdater {
 // CHECK-NORTC-NEXT:   GlobalMapUpdater() {
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[0], sycl::detail::kernel_args_sizes[0]);
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names[1], sycl::detail::kernel_args_sizes[1]);
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
 // CHECK-NORTC-NEXT:   }
 // CHECK-NORTC-NEXT: };
 // CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
index bef3bb9c35950..ac3cd76b2d92a 100644
--- a/sycl/include/sycl/detail/kernel_global_info.hpp
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -15,7 +15,8 @@ inline namespace _V1 {
 namespace detail {
 namespace free_function_info_map {
 
-__SYCL_EXPORT void add(const char *UniqueId, unsigned DeviceGlobalPtr);
+__SYCL_EXPORT void add(const char *const *UniqueId,
+                       const unsigned *DeviceGlobalPtr, unsigned Size);
 
 } // namespace free_function_info_map
 } // namespace detail
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index 860ec9ece1fed..05f0b250aa511 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -12,9 +12,14 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail::free_function_info_map {
 
-__SYCL_EXPORT void add(const char *UniqueId, unsigned KernelGlobalPtr) {
+__SYCL_EXPORT void add(const char *const *UniqueId,
+                       const unsigned *DeviceGlobalPtr, unsigned Size) {
+  std::unordered_map<std::string_view, unsigned> GlobalInfoToCopy;
+  for (size_t i = 0; i < Size; ++i) {
+    GlobalInfoToCopy[std::string_view{UniqueId[i]}] = DeviceGlobalPtr[i];
+  }
   detail::ProgramManager::getInstance().registerKernelGlobalInfo(
-      UniqueId, KernelGlobalPtr);
+      GlobalInfoToCopy);
 }
 
 } // namespace detail::free_function_info_map
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index ef4a634766ab5..190fb5bbc2316 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2407,13 +2407,21 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
 }
 
-void ProgramManager::registerKernelGlobalInfo(const char *UniqueId,
-                                              unsigned KernelGlobalPtr) {
-  m_KernelGlobalInfo.emplace(std::string_view(UniqueId), KernelGlobalPtr);
+void ProgramManager::registerKernelGlobalInfo(
+    std::unordered_map<std::string_view, unsigned> &GlobalInfoToCopy) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
+  if (m_KernelGlobalInfo.empty())
+    m_KernelGlobalInfo = std::move(GlobalInfoToCopy);
+  else {
+    for (auto &GlobalInfo : GlobalInfoToCopy) {
+      m_KernelGlobalInfo.insert(GlobalInfo);
+    }
+  }
 }
 
 std::optional<unsigned>
 ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
   const auto It = m_KernelGlobalInfo.find(UniqueId);
   if (It == m_KernelGlobalInfo.end())
     return std::nullopt;
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 029ed38bf88d9..cdfd7c905545a 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -257,7 +257,8 @@ class ProgramManager {
 
   // The function inserts or initializes a kernel global desc into the
   // kernel global map.
-  void registerKernelGlobalInfo(const char *UniqueId, unsigned KernelGlobalPtr);
+  void registerKernelGlobalInfo(
+      std::unordered_map<std::string_view, unsigned> &GlobalInfoToCopy);
 
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
@@ -548,7 +549,7 @@ class ProgramManager {
 
   // Maps between host_pipe identifiers and associated kernel global
   // information.
-  std::unordered_map<KernelNameStrT, unsigned> m_KernelGlobalInfo;
+  std::unordered_map<std::string_view, unsigned> m_KernelGlobalInfo;
 
   // Maps between host_pipe identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 43424e3eb1d82..9e003fcebf803 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3845,7 +3845,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
-_ZN4sycl3_V16detail22free_function_info_map3addEPKcj
+_ZN4sycl3_V16detail22free_function_info_map3addEPKPKcPKjj
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost10getSamplerEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost12getNumOfDimsEv

From 0b67f11e3a2acd4f94cb6e2c8fbdb055ce02206b Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Thu, 24 Jul 2025 13:24:20 +0200
Subject: [PATCH 21/23] [SYCL][TEST] update win abi test

---
 sycl/test/abi/sycl_symbols_windows.dump | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index d359b9e5c8dd2..fe412776c50b3 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -3765,7 +3765,7 @@
 ?__trunc_impl@_V1@sycl@@YANN@Z
 ?accelerator_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z
 ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
-?add@free_function_info_map@detail@_V1@sycl@@YAXPEBDPEBX@Z
+?add@free_function_info_map@detail@_V1@sycl@@YAXPEBQEBDPEBII@Z
 ?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA?AVnode@34567@AEBVproperty_list@67@@Z
 ?addAccessorReq@handler@_V1@sycl@@AEAAXV?$shared_ptr@VAccessorImplHost@detail@_V1@sycl@@@std@@@Z

From 532aff0d16f4c482d9ce03d0b40306da99936116 Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Thu, 24 Jul 2025 14:58:27 +0200
Subject: [PATCH 22/23] [SYCL] more informative names and comments

---
 sycl/source/detail/kernel_global_info.cpp            |  9 +++++----
 sycl/source/detail/kernel_impl.hpp                   |  2 ++
 .../detail/program_manager/program_manager.cpp       | 12 ++++++------
 .../detail/program_manager/program_manager.hpp       |  6 +++---
 4 files changed, 16 insertions(+), 13 deletions(-)

diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index 05f0b250aa511..fbc811ecf1227 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -12,14 +12,15 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail::free_function_info_map {
 
-__SYCL_EXPORT void add(const char *const *UniqueId,
-                       const unsigned *DeviceGlobalPtr, unsigned Size) {
+__SYCL_EXPORT void add(const char *const *FreeFunctionNames,
+                       const unsigned *FreeFunctionNumArgs, unsigned Size) {
   std::unordered_map<std::string_view, unsigned> GlobalInfoToCopy;
   for (size_t i = 0; i < Size; ++i) {
-    GlobalInfoToCopy[std::string_view{UniqueId[i]}] = DeviceGlobalPtr[i];
+    GlobalInfoToCopy[std::string_view{FreeFunctionNames[i]}] =
+        FreeFunctionNumArgs[i];
   }
   detail::ProgramManager::getInstance().registerKernelGlobalInfo(
-      GlobalInfoToCopy);
+      std::move(GlobalInfoToCopy));
 }
 
 } // namespace detail::free_function_info_map
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index 32c52309f1d80..4e74d5d4cfe5a 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -311,6 +311,8 @@ inline typename Param::return_type kernel_impl::get_info() const {
   static_assert(is_kernel_info_desc<Param>::value,
                 "Invalid kernel information descriptor");
   if constexpr (std::is_same_v<Param, info::kernel::num_args>) {
+    // if kernel is a free function, we need to get num_args from integration
+    // header, stored in program manager
     if (std::optional<unsigned> FFArgSize = getFreeFuncKernelArgSize())
       return *FFArgSize;
     checkIfValidForNumArgsInfoQuery();
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 190fb5bbc2316..8bdc499a7b769 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2408,13 +2408,13 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
 }
 
 void ProgramManager::registerKernelGlobalInfo(
-    std::unordered_map<std::string_view, unsigned> &GlobalInfoToCopy) {
+    std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy) {
   std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
-  if (m_KernelGlobalInfo.empty())
-    m_KernelGlobalInfo = std::move(GlobalInfoToCopy);
+  if (m_FreeFunctionKernelGlobalInfo.empty())
+    m_FreeFunctionKernelGlobalInfo = std::move(GlobalInfoToCopy);
   else {
     for (auto &GlobalInfo : GlobalInfoToCopy) {
-      m_KernelGlobalInfo.insert(GlobalInfo);
+      m_FreeFunctionKernelGlobalInfo.insert(GlobalInfo);
     }
   }
 }
@@ -2422,8 +2422,8 @@ void ProgramManager::registerKernelGlobalInfo(
 std::optional<unsigned>
 ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
   std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
-  const auto It = m_KernelGlobalInfo.find(UniqueId);
-  if (It == m_KernelGlobalInfo.end())
+  const auto It = m_FreeFunctionKernelGlobalInfo.find(UniqueId);
+  if (It == m_FreeFunctionKernelGlobalInfo.end())
     return std::nullopt;
   return It->second;
 }
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index cdfd7c905545a..cec86d7e1eaff 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -258,7 +258,7 @@ class ProgramManager {
   // The function inserts or initializes a kernel global desc into the
   // kernel global map.
   void registerKernelGlobalInfo(
-      std::unordered_map<std::string_view, unsigned> &GlobalInfoToCopy);
+      std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);
 
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
@@ -547,9 +547,9 @@ class ProgramManager {
   // their associated entry resources when they die.
   DeviceGlobalMap m_DeviceGlobals{/*OwnerControlledCleanup=*/true};
 
-  // Maps between host_pipe identifiers and associated kernel global
+  // Maps between free function kernel name and associated kernel global
   // information.
-  std::unordered_map<std::string_view, unsigned> m_KernelGlobalInfo;
+  std::unordered_map<std::string_view, unsigned> m_FreeFunctionKernelGlobalInfo;
 
   // Maps between host_pipe identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>

From be6f64c4b891054b37ef1679006db1f1d858f4eb Mon Sep 17 00:00:00 2001
From: "Klochkov, Denis" <denis.klochkov@intel.com>
Date: Tue, 29 Jul 2025 01:47:25 -0700
Subject: [PATCH 23/23] [SYCL] update formatting

---
 clang/lib/Sema/SemaSYCL.cpp                                     | 2 +-
 .../CodeGenSYCL/free_function_default_template_arguments.cpp    | 2 +-
 clang/test/CodeGenSYCL/free_function_int_header.cpp             | 2 +-
 clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp    | 2 +-
 4 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 732d36904d24c..f983582ed2085 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7196,7 +7196,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "namespace detail {\n";
     O << "struct GlobalMapUpdater {\n";
     O << "  GlobalMapUpdater() {\n";
-    O << "sycl::detail::free_function_info_map::add("
+    O << "    sycl::detail::free_function_info_map::add("
       << "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
       << KernelDescs.size() << ");\n";
     O << "  }\n";
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index dc542cb50f389..2a9187d618b24 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -1092,7 +1092,7 @@ namespace Testing::Tests {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:   GlobalMapUpdater() {
-// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
+// CHECK-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index add97822677f4..b05c299a2e478 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1572,7 +1572,7 @@ void ff_24(int arg) {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:  GlobalMapUpdater() {
-// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
+// CHECK-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
 // CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 0009e71c7c9f1..48d5662e709aa 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -81,7 +81,7 @@ int main(){
 // CHECK-NORTC-NEXT: namespace detail {
 // CHECK-NORTC-NEXT: struct GlobalMapUpdater {
 // CHECK-NORTC-NEXT:   GlobalMapUpdater() {
-// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes,
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
 // CHECK-NORTC-NEXT:   }
 // CHECK-NORTC-NEXT: };
 // CHECK-NORTC-NEXT: static GlobalMapUpdater updater;