diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 07a6c25dced91..f983582ed2085 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6666,6 +6666,34 @@ 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 << "\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 << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr unsigned getNumParams() { return " << KParamsSize
+      << "; }\n";
+    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";
+    O << "\n";
+  }
+
 private:
   /// Helper method to get string with template types
   /// \param TAL The template argument list.
@@ -6915,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";
@@ -7127,6 +7160,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;
 
@@ -7156,22 +7190,21 @@ 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 << "#include <sycl/detail/kernel_global_info.hpp>\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{\""
-      << K.Name << "\"});\n";
-    O << "}\n";
-    O << "}\n";
-    ++ShimCounter;
+    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, 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";
   }
 }
 
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 2debb64fcc8a3..2a9187d618b24 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;
@@ -665,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-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;
+// 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 48a03c6c65916..b05c299a2e478 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;
@@ -1077,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{"_Z18__sycl_kernel_ff_2Piii"});
-// 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{"_Z18__sycl_kernel_ff_2Piiii"});
-// 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{"_Z18__sycl_kernel_ff_3IiEvPT_S0_S0_"});
-// 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{"_Z18__sycl_kernel_ff_3IfEvPT_S0_S0_"});
+// CHECK-NEXT: inline namespace _V1 {
+// 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, 33);
 // 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{"_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"});
-// 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{"_Z18__sycl_kernel_ff_410NoPointers8Pointers3Agg"});
-// 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{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
-// 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{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
-// 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{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"});
-// 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{"_ZN28__sycl_kernel_free_functions4ff_9EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5tests5ff_10EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi"});
-// 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{"_ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5ff_13EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5tests5ff_13EiPi"});
-// 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{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"});
-
-// 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: 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: 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: 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: 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: 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: }
-// 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{"_ZN28__sycl_kernel_free_functions5tests5ff_14EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5ff_15EiPi"});
-// 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{"_ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_"});
-// 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{"_ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_"});
-// 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{"_ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_"});
-// 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{"_Z19__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE"});
-// 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: }
-// 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: // 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: }
-// 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{"_Z19__sycl_kernel_ff_23i"});
-
-// 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 29b697691f445..48d5662e709aa 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -73,8 +73,18 @@ 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_kernel_free_function_singlePiii"});
 
-// 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: #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(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
+// 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/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/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
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/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
new file mode 100644
index 0000000000000..ac3cd76b2d92a
--- /dev/null
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -0,0 +1,24 @@
+//==-------------------- 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 char *const *UniqueId,
+                       const unsigned *DeviceGlobalPtr, unsigned Size);
+
+} // namespace free_function_info_map
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
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_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index e2709a94e4be3..d8a7a0cd9cd99 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -732,6 +732,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/CMakeLists.txt b/sycl/source/CMakeLists.txt
index c564c63b06b3a..24a471eacb0f6 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 15a23983e5c54..839d0b0916da9 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -1013,6 +1013,12 @@ class kernel_bundle_impl
 
   DeviceGlobalMap &getDeviceGlobalMap() { return MDeviceGlobals; }
 
+  std::optional<unsigned>
+  tryGetKernelArgsSize(const std::string_view KernelName) const {
+    auto &PM = sycl::detail::ProgramManager::getInstance();
+    return PM.getKernelGlobalInfoDesc(KernelName.data());
+  }
+
 private:
   DeviceGlobalMapEntry *getDeviceGlobalEntry(const std::string &Name) const {
     if (!hasSourceBasedImages() && !hasSYCLBINImages()) {
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
new file mode 100644
index 0000000000000..fbc811ecf1227
--- /dev/null
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -0,0 +1,28 @@
+//==-------------------- 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 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{FreeFunctionNames[i]}] =
+        FreeFunctionNumArgs[i];
+  }
+  detail::ProgramManager::getInstance().registerKernelGlobalInfo(
+      std::move(GlobalInfoToCopy));
+}
+
+} // namespace detail::free_function_info_map
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 0616a4ec71e5c..cb78b8f50e1c5 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -138,6 +138,10 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
       "interoperability function or to query a device built-in kernel");
 }
 
+std::optional<unsigned> kernel_impl ::getFreeFuncKernelArgSize() const {
+  return MKernelBundleImpl->tryGetKernelArgsSize(getName());
+}
+
 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 bd4d09c46a686..c6135f1fea9d9 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -266,6 +266,7 @@ class kernel_impl {
                                size_t DynamicLocalMemorySize) const;
 
   void enableUSMIndirectAccess() const;
+  std::optional<unsigned> getFreeFuncKernelArgSize() const;
 };
 
 template <int Dimensions>
@@ -305,9 +306,13 @@ 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>) {
+    // 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();
-
+  }
   return get_kernel_info<Param>(this->getHandleRef(), getAdapter());
 }
 
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 2bdca5ad35acc..62e8befe00c73 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2378,6 +2378,27 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
 }
 
+void ProgramManager::registerKernelGlobalInfo(
+    std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
+  if (m_FreeFunctionKernelGlobalInfo.empty())
+    m_FreeFunctionKernelGlobalInfo = std::move(GlobalInfoToCopy);
+  else {
+    for (auto &GlobalInfo : GlobalInfoToCopy) {
+      m_FreeFunctionKernelGlobalInfo.insert(GlobalInfo);
+    }
+  }
+}
+
+std::optional<unsigned>
+ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
+  const auto It = m_FreeFunctionKernelGlobalInfo.find(UniqueId);
+  if (It == m_FreeFunctionKernelGlobalInfo.end())
+    return std::nullopt;
+  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 42ac0475d90f7..87ce40d1a7e98 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -255,6 +255,15 @@ 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 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.
+  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);
 
@@ -535,6 +544,10 @@ class ProgramManager {
   // their associated entry resources when they die.
   DeviceGlobalMap m_DeviceGlobals{/*OwnerControlledCleanup=*/true};
 
+  // Maps between free function kernel name and associated kernel global
+  // information.
+  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>>
       m_HostPipes;
diff --git a/sycl/test-e2e/FreeFunctionKernels/num_args.cpp b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
new file mode 100644
index 0000000000000..701b900036a62
--- /dev/null
+++ b/sycl/test-e2e/FreeFunctionKernels/num_args.cpp
@@ -0,0 +1,86 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.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, 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>();
+  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;
+}
+
+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();
+  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, 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 dd392cf315b88..71269e088a88a 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3849,6 +3849,7 @@ _ZNK4sycl3_V16detail19kernel_bundle_plain5emptyEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail21LocalAccessorBaseHost7getSizeEv
+_ZN4sycl3_V16detail22free_function_info_map3addEPKPKcPKjj
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost10getSamplerEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost11getPropListEv
 _ZNK4sycl3_V16detail28SampledImageAccessorBaseHost12getNumOfDimsEv
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index 748d74482a1a5..195b20a66d934 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -3768,6 +3768,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@@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