diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 861ac6241e308..079ef4c3915ea 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -15,10 +15,9 @@
 #include "clang/AST/QualTypeNames.h"
 #include "clang/AST/RecordLayout.h"
 #include "clang/AST/RecursiveASTVisitor.h"
-#include "clang/AST/TemplateArgumentVisitor.h"
-#include "clang/AST/Mangle.h"
 #include "clang/AST/SYCLKernelInfo.h"
 #include "clang/AST/StmtSYCL.h"
+#include "clang/AST/TemplateArgumentVisitor.h"
 #include "clang/AST/TypeOrdering.h"
 #include "clang/AST/TypeVisitor.h"
 #include "clang/Analysis/CallGraph.h"
@@ -27,7 +26,6 @@
 #include "clang/Basic/Diagnostic.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/Version.h"
-#include "clang/AST/SYCLKernelInfo.h"
 #include "clang/Sema/Attr.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/ParsedAttr.h"
@@ -6425,6 +6423,120 @@ static void EmitPragmaDiagnosticPop(raw_ostream &O) {
   O << "\n";
 }
 
+template <typename BeforeFn, typename AfterFn>
+static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS,
+                          const DeclContext *DC) {
+  if (DC->isTranslationUnit())
+    return;
+
+  const auto *CurDecl = cast<Decl>(DC);
+  // Ensure we are in the canonical version, so that we know we have the 'full'
+  // name of the thing.
+  CurDecl = CurDecl->getCanonicalDecl();
+
+  // We are intentionally skipping linkage decls and record decls.  Namespaces
+  // can appear in a linkage decl, but not a record decl, so we don't have to
+  // worry about the names getting messed up from that.  We handle record decls
+  // later when printing the name of the thing.
+  const auto *NS = dyn_cast<NamespaceDecl>(CurDecl);
+  if (NS)
+    Before(OS, NS);
+
+  if (const DeclContext *NewDC = CurDecl->getDeclContext())
+    PrintNSHelper(Before, After, OS, NewDC);
+
+  if (NS)
+    After(OS, NS);
+}
+
+static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC,
+                            bool isPrintNamesOnly = false) {
+  PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {},
+                [isPrintNamesOnly](raw_ostream &OS, const NamespaceDecl *NS) {
+                  if (!isPrintNamesOnly) {
+                    if (NS->isInline())
+                      OS << "inline ";
+                    OS << "namespace ";
+                  }
+                  if (!NS->isAnonymousNamespace()) {
+                    OS << NS->getName();
+                    if (isPrintNamesOnly)
+                      OS << "::";
+                    else
+                      OS << " ";
+                  }
+                  if (!isPrintNamesOnly) {
+                    OS << "{\n";
+                  }
+                },
+                OS, DC);
+}
+
+static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
+  PrintNSHelper(
+      [](raw_ostream &OS, const NamespaceDecl *NS) {
+        OS << "} // ";
+        if (NS->isInline())
+          OS << "inline ";
+
+        OS << "namespace ";
+        if (!NS->isAnonymousNamespace())
+          OS << NS->getName();
+
+        OS << '\n';
+      },
+      [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC);
+}
+
+class FreeFunctionPrinter {
+  raw_ostream &O;
+  const PrintingPolicy &Policy;
+  bool NSInserted = false;
+
+public:
+  FreeFunctionPrinter(raw_ostream &O, const PrintingPolicy &Policy)
+      : O(O), Policy(Policy) {}
+
+  /// Emits the function declaration of a free function.
+  /// \param FD The function declaration to print.
+  /// \param Args The arguments of the function.
+  void printFreeFunctionDeclaration(const FunctionDecl *FD,
+                                    const std::string &Args) {
+    const DeclContext *DC = FD->getDeclContext();
+    if (DC) {
+      // if function in namespace, print namespace
+      if (isa<NamespaceDecl>(DC)) {
+        PrintNamespaces(O, FD);
+        // Set flag to print closing braces for namespaces and namespace in shim
+        // function
+        NSInserted = true;
+      }
+      O << FD->getReturnType().getAsString() << " ";
+      O << FD->getNameAsString() << "(" << Args << ");";
+      if (NSInserted) {
+        O << "\n";
+        PrintNSClosingBraces(O, FD);
+      }
+      O << "\n";
+    }
+  }
+
+  /// Emits free function shim function.
+  /// \param FD The function declaration to print.
+  /// \param ShimCounter The counter for the shim function.
+  /// \param ParmList The parameter list of the function.
+  void printFreeFunctionShim(const FunctionDecl *FD, const unsigned ShimCounter,
+                             const std::string &ParmList) {
+    // Generate a shim function that returns the address of the free function.
+    O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n";
+    O << "  return (void (*)(" << ParmList << "))";
+
+    if (NSInserted)
+      PrintNamespaces(O, FD, /*isPrintNamesOnly=*/true);
+    O << FD->getIdentifier()->getName().data();
+  }
+};
+
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
@@ -6713,16 +6825,25 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     if (K.SyclKernel->getLanguageLinkage() == CLanguageLinkage)
       O << "extern \"C\" ";
     std::string ParmList;
+    std::string ParmListWithNames;
     bool FirstParam = true;
     Policy.SuppressDefaultTemplateArgs = false;
     Policy.PrintCanonicalTypes = true;
+    llvm::raw_string_ostream ParmListWithNamesOstream{ParmListWithNames};
     for (ParmVarDecl *Param : K.SyclKernel->parameters()) {
       if (FirstParam)
         FirstParam = false;
-      else
+      else {
         ParmList += ", ";
+        ParmListWithNamesOstream << ", ";
+      }
+      Policy.SuppressTagKeyword = true;
+      Param->getType().print(ParmListWithNamesOstream, Policy);
+      Policy.SuppressTagKeyword = false;
+      ParmListWithNamesOstream << " " << Param->getNameAsString();
       ParmList += Param->getType().getCanonicalType().getAsString(Policy);
     }
+    ParmListWithNamesOstream.flush();
     FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate();
     Policy.PrintCanonicalTypes = false;
     Policy.SuppressDefinition = true;
@@ -6756,17 +6877,15 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     // template arguments that match default template arguments while printing
     // template-ids, even if the source code doesn't reference them.
     Policy.EnforceDefaultTemplateArgs = true;
+    FreeFunctionPrinter FFPrinter(O, Policy);
     if (FTD) {
       FTD->print(O, Policy);
+      O << ";\n";
     } else {
-      K.SyclKernel->print(O, Policy);
+      FFPrinter.printFreeFunctionDeclaration(K.SyclKernel, ParmListWithNames);
     }
-    O << ";\n";
 
-    // Generate a shim function that returns the address of the free function.
-    O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n";
-    O << "  return (void (*)(" << ParmList << "))"
-      << K.SyclKernel->getIdentifier()->getName().data();
+    FFPrinter.printFreeFunctionShim(K.SyclKernel, ShimCounter, ParmList);
     if (FTD) {
       const TemplateArgumentList *TAL =
           K.SyclKernel->getTemplateSpecializationArgs();
@@ -6935,61 +7054,6 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
   return emit(Out);
 }
 
-template <typename BeforeFn, typename AfterFn>
-static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS,
-                          const DeclContext *DC) {
-  if (DC->isTranslationUnit())
-    return;
-
-  const auto *CurDecl = cast<Decl>(DC);
-  // Ensure we are in the canonical version, so that we know we have the 'full'
-  // name of the thing.
-  CurDecl = CurDecl->getCanonicalDecl();
-
-  // We are intentionally skipping linkage decls and record decls.  Namespaces
-  // can appear in a linkage decl, but not a record decl, so we don't have to
-  // worry about the names getting messed up from that.  We handle record decls
-  // later when printing the name of the thing.
-  const auto *NS = dyn_cast<NamespaceDecl>(CurDecl);
-  if (NS)
-    Before(OS, NS);
-
-  if (const DeclContext *NewDC = CurDecl->getDeclContext())
-    PrintNSHelper(Before, After, OS, NewDC);
-
-  if (NS)
-    After(OS, NS);
-}
-
-static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC) {
-  PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {},
-                [](raw_ostream &OS, const NamespaceDecl *NS) {
-                  if (NS->isInline())
-                    OS << "inline ";
-                  OS << "namespace ";
-                  if (!NS->isAnonymousNamespace())
-                    OS << NS->getName() << " ";
-                  OS << "{\n";
-                },
-                OS, DC);
-}
-
-static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
-  PrintNSHelper(
-      [](raw_ostream &OS, const NamespaceDecl *NS) {
-        OS << "} // ";
-        if (NS->isInline())
-          OS << "inline ";
-
-        OS << "namespace ";
-        if (!NS->isAnonymousNamespace())
-          OS << NS->getName();
-
-        OS << '\n';
-      },
-      [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC);
-}
-
 static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter,
                             const std::string &LastShim,
                             const NamespaceDecl *AnonNS) {
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 62a121d218b8b..f7da6e8f1f772 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -86,18 +86,20 @@ foo(Arg1<int> arg) {
 // CHECK-NEXT: template <typename T, typename, int a, typename, typename ...TS> struct Arg;
 // CHECK-NEXT: }
 
-// CHECK: void ns::simple(ns::Arg<char, int, 12, ns::notatuple>);
-// CHECK-NEXT: static constexpr auto __sycl_shim1() {
-// CHECK-NEXT:   return (void (*)(struct ns::Arg<char, int, 12, struct ns::notatuple>))simple;
+// CHECK: namespace ns {
+// CHECK-NEXT: void simple(ns::Arg<char, int, 12, ns::notatuple> );
+// CHECK-NEXT: } // namespace ns
+// CHECK: static constexpr auto __sycl_shim1() {
+// CHECK-NEXT:   return (void (*)(struct ns::Arg<char, int, 12, struct ns::notatuple>))ns::simple;
 // CHECK-NEXT: }
 
 // CHECK: Forward declarations of kernel and its argument types:
 // CHECK: namespace ns {
 // CHECK: namespace ns1 {
 // CHECK-NEXT: template <typename A> class hasDefaultArg;
-// CHECK-NEXT: }
+// CHECK-NEXT: }}
 
-// CHECK: void simple1(ns::Arg<ns::ns1::hasDefaultArg<ns::notatuple>, int, 12, ns::notatuple>);
+// CHECK: void simple1(ns::Arg<ns::ns1::hasDefaultArg<ns::notatuple>, int, 12, ns::notatuple> );
 // CHECK-NEXT: static constexpr auto __sycl_shim2() {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, 12, struct ns::notatuple>))simple1;
 // CHECK-NEXT: }
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index 9e5fdd9fd495f..4fe57f8acf34b 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -78,6 +78,7 @@ __attribute__((sycl_device))
 template void ff_6(Agg S1, Derived S2, int);
 
 constexpr int TestArrSize = 3;
+constexpr int TestArrSizeAlias = 50;
 
 template <int ArrSize>
 struct KArgWithPtrArray {
@@ -87,6 +88,18 @@ struct KArgWithPtrArray {
   constexpr int getArrSize() { return ArrSize; }
 };
 
+namespace free_functions {
+  template <int ArrSize>
+  struct KArgWithPtrArray {
+    float *data[ArrSize];
+    float start[ArrSize];
+    float end[ArrSize];
+    constexpr int getArrSize() { return ArrSize; }
+  };
+
+  using AliasStruct = KArgWithPtrArray<TestArrSizeAlias>;
+}
+
 template <int ArrSize>
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
 void ff_7(KArgWithPtrArray<ArrSize> KArg) {
@@ -102,11 +115,106 @@ __attribute__((sycl_device))
 void ff_8(sycl::work_group_memory<int>) {
 }
 
+// function in namespace
+namespace free_functions {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_9(int start, int *ptr) {
+}
+}
+
+// function in nested namespace
+namespace free_functions::tests {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_10(int start, int *ptr) {
+}
+}
+
+// function in inline namespace
+namespace free_functions::tests {
+inline namespace V1 {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_11(int start, int *ptr) {
+}
+}
+}
+
+//function in anonymous namespace
+namespace {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_12(int start, int *ptr) {
+}
+}
+
+// functions with the same name but in different namespaces
+namespace free_functions {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_13(int start, int *ptr) {
+}
+}
+namespace free_functions::tests {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_13(int start, int *ptr) {
+}
+}
+
 __attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] 
 void ff_9(sycl::dynamic_work_group_memory<int>) {
 }
 
+typedef int TypedefType;
+using AliasType = Derived;
+
+namespace free_functions::tests {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_14(TypedefType start, TypedefType *ptr) {
+}
+}
+
+namespace free_functions::tests {
+typedef int NamespaceTypedefType;
+using AliasType = Agg;
+}
+
+namespace free_functions {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_15(free_functions::tests::NamespaceTypedefType start, free_functions::tests::NamespaceTypedefType *ptr) {
+}
+}
+
+namespace free_functions {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_16(free_functions::tests::AliasType start, free_functions::tests::AliasType *ptr) {
+}
+}
+
+namespace free_functions {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_17(AliasType start, AliasType *ptr) {
+}
+}
+
+namespace free_functions {
+  struct Agg {
+    int a;
+    float b;
+  };
+}
+
+namespace free_functions::tests {
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_18(free_functions::Agg start, free_functions::Agg *ptr) {
+  ptr->a = start.a + 1;
+  ptr->b = start.b + 1.1f;
+}
+}
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_19(free_functions::AliasStruct KArg) {
+  for (int j = 0; j < TestArrSizeAlias; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
 
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
@@ -118,7 +226,23 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
+
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions4ff_9EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests5ff_10EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests2V15ff_11EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel__GLOBAL__N_15ff_12EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5ff_13EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests5ff_13EiPi
+
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
+
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests5ff_14EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5ff_15EiPi
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5ff_16E3AggPS0_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5ff_17E7DerivedPS0_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE
+
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
@@ -165,15 +289,63 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 // CHECK:  //--- _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
 // CHECK-NEXT:  { kernel_param_kind_t::kind_work_group_memory, 8, 0 },
 
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions4ff_9EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5ff_13EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 }, 
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
 // CHECK:  //--- _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
 // CHECK-NEXT:  { kernel_param_kind_t::kind_dynamic_work_group_memory, 8, 0 },
 
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5ff_15EiPi
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 4 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 32, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 32 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 40, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 40 },
+
+// CHECK:  //--- _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 8, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_pointer, 8, 8 },
+
+// CHECK: //--- _Z19__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 800, 0 },
+
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
 // CHECK-NEXT: };
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_2Piii as a free function kernel
 // CHECK: Forward declarations of kernel and its argument types:
-// CHECK: void ff_2(int *ptr, int start, int end);
+// CHECK: void ff_2(int * ptr, int start, int end);
 // CHECK-NEXT: static constexpr auto __sycl_shim1() {
 // CHECK-NEXT:   return (void (*)(int *, int, int))ff_2;
 // CHECK-NEXT: }
@@ -190,7 +362,7 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 
 // CHECK: Definition of _Z18__sycl_kernel_ff_2Piiii as a free function kernel
 // CHECK: Forward declarations of kernel and its argument types:
-// CHECK: void ff_2(int *ptr, int start, int end, int value);
+// CHECK: void ff_2(int * ptr, int start, int end, int value);
 // CHECK-NEXT: static constexpr auto __sycl_shim2() {
 // CHECK-NEXT:   return (void (*)(int *, int, int, int))ff_2;
 // CHECK-NEXT: }
@@ -319,7 +491,7 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 // CHECK: Forward declarations of kernel and its argument types:
 // CHECK: template <typename DataT> class work_group_memory;
 
-// CHECK: void ff_8(sycl::work_group_memory<int>);
+// CHECK: void ff_8(sycl::work_group_memory<int> );
 // CHECK-NEXT: static constexpr auto __sycl_shim9() {
 // CHECK-NEXT: return (void (*)(class sycl::work_group_memory<int>))ff_8;
 // CHECK-NEXT: }
@@ -334,26 +506,272 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 // CHECK-NEXT: };
 // CHECK-NEXT: }
 
+
+// CHECK: Definition of _ZN28__sycl_kernel_free_functions4ff_9EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: namespace free_functions {
+// CHECK-NEXT: void ff_9(int start, int * ptr);
+// CHECK-NEXT: } // namespace free_functions
+
+// CHECK: static constexpr auto __sycl_shim10() {
+// CHECK-NEXT:   return (void (*)(int, int *))free_functions::ff_9;
+// CHECK-NEXT: }
+
+// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+
+// CHECK: namespace free_functions {
+// CHECK-NEXT: namespace tests {
+// CHECK-NEXT: void ff_10(int start, int * ptr);
+// CHECK-NEXT: } // namespace tests
+// CHECK-NEXT: } // namespace free_functions
+// CHECK: static constexpr auto __sycl_shim11() {
+// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_10;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim11()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim11()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
+// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  namespace tests {
+// CHECK-NEXT:  inline namespace V1 {
+// CHECK-NEXT:  void ff_11(int start, int * ptr);
+// CHECK-NEXT:  } // inline namespace V1
+// CHECK-NEXT:  } // namespace tests
+// CHECK-NEXT:  } // namespace free_functions
+// 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-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim12()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim12()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: Definition of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+
+// CHECK: namespace {
+// CHECK-NEXT:  void ff_12(int start, int * ptr);
+// CHECK-NEXT:  } // namespace 
+// CHECK:  static constexpr auto __sycl_shim13() {
+// CHECK-NEXT:    return (void (*)(int, int *))ff_12;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim13()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim13()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: Definition of _ZN28__sycl_kernel_free_functions5ff_13EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  void ff_13(int start, int * ptr);
+// CHECK-NEXT:  } // namespace free_functions
+// CHECK:  static constexpr auto __sycl_shim14() {
+// CHECK-NEXT:    return (void (*)(int, int *))free_functions::ff_13;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim14()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim14()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+ 
+// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi as a free function kernel
+// CHECK: Forward declarations of kernel and its argument types:
+
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  namespace tests {
+// CHECK-NEXT:  void ff_13(int start, int * ptr);
+// CHECK-NEXT:  } // namespace tests
+// CHECK-NEXT:  } // namespace free_functions
+// CHECK:  static constexpr auto __sycl_shim15() {
+// CHECK-NEXT:    return (void (*)(int, int *))free_functions::tests::ff_13;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim15()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim15()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+
 // CHECK:  // Definition of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE as a free function kernel 
-//
 // CHECK: Forward declarations of kernel and its argument types:
-// CHECK: template <typename DataT> class dynamic_work_group_memory;
+// CHECK-NEXT: namespace sycl { inline namespace _V1 {
+// CHECK-NEXT: template <typename DataT> class dynamic_work_group_memory;
+// CHECK-NEXT: }}
 
-// CHECK: void ff_9(sycl::dynamic_work_group_memory<int>);
-// CHECK-NEXT: static constexpr auto __sycl_shim10() {
+// CHECK: void ff_9(sycl::dynamic_work_group_memory<int> );
+// 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-NEXT: template <>
-// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> {
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim16()> {
 // CHECK-NEXT: static constexpr bool value = true;
 // CHECK-NEXT: };
 // CHECK-NEXT: template <>
-// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> {
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim16()> {
 // CHECK-NEXT: static constexpr bool value = true;
 // CHECK-NEXT: };
 // CHECK-NEXT: }
 
+// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  namespace tests {
+// CHECK-NEXT:  void ff_14(int start, int * ptr);
+// CHECK-NEXT:  } // namespace tests
+// CHECK-NEXT:  } // namespace free_functions
+  
+// CHECK:  static constexpr auto __sycl_shim17() {
+// CHECK-NEXT:    return (void (*)(int, int *))free_functions::tests::ff_14;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim17()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim17()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_15EiPi as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  void ff_15(int start, int * ptr);
+// CHECK-NEXT:  } // namespace free_functions
+// CHECK:  static constexpr auto __sycl_shim18() {
+// CHECK-NEXT:    return (void (*)(int, int *))free_functions::ff_15;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim18()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim18()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_ as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  void ff_16(Agg start, Agg * ptr);
+// CHECK-NEXT:  } // namespace free_functions
+// CHECK:  static constexpr auto __sycl_shim19() {
+// CHECK-NEXT:    return (void (*)(struct Agg, struct Agg *))free_functions::ff_16;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim19()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim19()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_ as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK: namespace free_functions {
+// CHECK-NEXT:  void ff_17(Derived start, Derived * ptr);
+// CHECK-NEXT:  } // namespace free_functions
+// CHECK:  static constexpr auto __sycl_shim20() {
+// CHECK-NEXT:    return (void (*)(struct Derived, struct Derived *))free_functions::ff_17;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim20()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim20()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
+// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_ as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK-NEXT: namespace free_functions { 
+// CHECK-NEXT:   struct Agg;
+// CHECK-NEXT:   }
+// CHECK:   namespace free_functions {
+// CHECK-NEXT:   namespace tests {
+// CHECK-NEXT:   void ff_18(free_functions::Agg start, free_functions::Agg * ptr);
+// CHECK-NEXT:   } // namespace tests
+// CHECK-NEXT:   } // namespace free_functions
+// CHECK:   static constexpr auto __sycl_shim21() {
+// CHECK-NEXT:     return (void (*)(struct free_functions::Agg, struct free_functions::Agg *))free_functions::tests::ff_18;
+// CHECK-NEXT:   }
+// CHECK-NEXT:   namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim21()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim21()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   }
+
+// CHECK: // Definition of _Z19__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE as a free function kernel
+// CHECK: // Forward declarations of kernel and its argument types:
+// CHECK-NEXT: namespace free_functions { 
+// CHECK-NEXT:  template <int ArrSize> struct KArgWithPtrArray;
+// CHECK-NEXT:  }
+  
+// CHECK:  void ff_19(free_functions::KArgWithPtrArray<50> KArg);
+// CHECK-NEXT:  static constexpr auto __sycl_shim22() {
+// CHECK-NEXT:    return (void (*)(struct free_functions::KArgWithPtrArray<50>))ff_19;
+// CHECK-NEXT:  }
+// CHECK-NEXT:  namespace sycl {
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_kernel<__sycl_shim22()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  template <>
+// CHECK-NEXT:  struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim22()> {
+// CHECK-NEXT:    static constexpr bool value = true;
+// CHECK-NEXT:  };
+// CHECK-NEXT:  }
+
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
 // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -427,11 +845,101 @@ void ff_9(sycl::dynamic_work_group_memory<int>) {
 // 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_shim10()>() {
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim16()>() {
 // CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"});
 // CHECK-NEXT: }
 // 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_shim17()>() {
+// 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_shim18()>() {
+// 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_shim19()>() {
+// 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_shim20()>() {
+// 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_shim21()>() {
+// CHECK-NEXT:    return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_"});
+// CHECK-NEXT:  }
+// CHECK-NEXT:  }
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index 214318b563fa8..29b697691f445 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -52,7 +52,7 @@ int main(){
 // CHECK-RTC-NOT: free_function_single_kernel
 // CHECK-RTC-NOT: free_function_nd_range
 
-// CHECK-NORTC:       void free_function_single(int *ptr, int start, int end);
+// CHECK-NORTC:       void free_function_single(int * ptr, int start, int end);
 // CHECK-NORTC:       static constexpr auto __sycl_shim[[#FIRST:]]()
 // CHECK-NORTC-NEXT:  return (void (*)(int *, int, int))free_function_single;
 
@@ -63,7 +63,7 @@ int main(){
 // CHECK-NORTC-NEXT:  static constexpr bool value = true;
 
 
-// CHECK-NORTC:       void free_function_nd_range(int *ptr, int start, int end);
+// CHECK-NORTC:       void free_function_nd_range(int * ptr, int start, int end);
 // CHECK-NORTC:       static constexpr auto __sycl_shim[[#SECOND:]]() {
 // CHECK-NORTC-NEXT:  return (void (*)(int *, int, int))free_function_nd_range;
 
diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp
new file mode 100644
index 0000000000000..472f7bdb15f78
--- /dev/null
+++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp
@@ -0,0 +1,235 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} %cxx_std_optionc++20 -o %t.out
+// RUN: %{run} %t.out
+
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+#include <type_traits>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void func(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+
+namespace free_functions::tests {
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_in_ns(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void func(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+} // namespace free_functions::tests
+
+namespace free_functions::tests {
+inline namespace V1 {
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_in_inline_ns(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+} // namespace V1
+} // namespace free_functions::tests
+
+namespace {
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_in_anonymous_ns(float start, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = start + static_cast<float>(id);
+}
+} // namespace
+
+struct TestClass {
+  float data;
+  TestClass(float d) : data(d) {}
+};
+
+template <typename T> struct TemplatedTestClass {
+  T data;
+  TemplatedTestClass(T d) : data(d) {}
+};
+
+using IntClassAlias = TemplatedTestClass<int>;
+using FloatClassAlias = TemplatedTestClass<float>;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_with_test_class(float start, TestClass *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id].data = start + static_cast<float>(id);
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_with_int_alias_test_class(float start, IntClassAlias *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id].data = start + static_cast<int>(id);
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void function_with_float_alias_test_class(float start, FloatClassAlias *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id].data = start + static_cast<float>(id);
+}
+
+template <typename T>
+concept NumericType = std::is_arithmetic_v<std::remove_reference_t<T>>;
+
+template <typename T>
+  requires NumericType<T>
+void check_result(T *ptr) {
+  for (size_t i = 0; i < NUM; ++i) {
+    const float expected = 3.14f + static_cast<float>(i);
+    assert(ptr[i] == expected &&
+           "Kernel execution did not produce the expected result");
+  }
+}
+
+template <typename T>
+concept HasDataMemeber = requires(T t) {
+  { t.data } -> NumericType;
+};
+
+template <typename T>
+  requires HasDataMemeber<T>
+void check_result(T *ptr) {
+  using DataType = decltype(ptr->data);
+  for (size_t i = 0; i < NUM; ++i) {
+    const DataType expected = 3.14f + static_cast<DataType>(i);
+    assert(ptr[i].data == expected &&
+           "Kernel execution did not produce the expected result");
+  }
+}
+
+template <typename T>
+static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) {
+  T *ptr = sycl::malloc_shared<T>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(3.14f, ptr);
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, kernel);
+   }).wait();
+  check_result<T>(ptr);
+}
+
+void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "func".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<func, sycl::bundle_state::executable>(ctxt);
+  // Get a kernel object for the "func" function from that bundle.
+  sycl::kernel k_func = exe_bndl.ext_oneapi_get_kernel<func>();
+  call_kernel_code<float>(q, k_func);
+}
+
+void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "function_in_ns".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<free_functions::tests::function_in_ns,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "function_in_ns" function from that bundle.
+  sycl::kernel k_function_in_ns =
+      exe_bndl.ext_oneapi_get_kernel<free_functions::tests::function_in_ns>();
+  call_kernel_code<float>(q, k_function_in_ns);
+}
+
+void test_func_in_ns_with_same_name(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "func".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<free_functions::tests::func,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "func" function from that bundle.
+  sycl::kernel k_func_in_ns =
+      exe_bndl.ext_oneapi_get_kernel<free_functions::tests::func>();
+  call_kernel_code<float>(q, k_func_in_ns);
+}
+
+void test_function_in_inline_ns(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "function_in_inline_ns".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<free_functions::tests::function_in_inline_ns,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "function_in_inline_ns" function from that
+  // bundle.
+  sycl::kernel k_function_in_inline_ns = exe_bndl.ext_oneapi_get_kernel<
+      free_functions::tests::function_in_inline_ns>();
+  call_kernel_code<float>(q, k_function_in_inline_ns);
+}
+
+void test_function_in_anonymous_ns(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "function_in_anonymous_ns".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<function_in_anonymous_ns,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "function_in_anonymous_ns" function from that
+  // bundle.
+  sycl::kernel k_function_in_anonymous_ns =
+      exe_bndl.ext_oneapi_get_kernel<function_in_anonymous_ns>();
+  call_kernel_code<float>(q, k_function_in_anonymous_ns);
+}
+
+void test_function_with_class(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "function_with_test_class".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<function_with_test_class,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "function_with_test_class" function from that
+  // bundle.
+  sycl::kernel k_function_with_test_class =
+      exe_bndl.template ext_oneapi_get_kernel<function_with_test_class>();
+  call_kernel_code<TestClass>(q, k_function_with_test_class);
+}
+
+void test_fucntions_with_int_class_alias(sycl::queue &q, sycl::context &ctxt) {
+  // Get a kernel bundle that contains the free function kernel
+  // "function_with_int_alias_test_class".
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<function_with_int_alias_test_class,
+                                 sycl::bundle_state::executable>(ctxt);
+
+  // Get a kernel object for the "function_with_int_alias_test_class" function
+  // from that bundle.
+  sycl::kernel k_function_with_int_alias_test_class =
+      exe_bndl
+          .template ext_oneapi_get_kernel<function_with_int_alias_test_class>();
+  call_kernel_code<IntClassAlias>(q, k_function_with_int_alias_test_class);
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctxt = q.get_context();
+
+  test_function_without_ns(q, ctxt);
+  test_function_in_ns(q, ctxt);
+  test_function_in_inline_ns(q, ctxt);
+  test_function_in_anonymous_ns(q, ctxt);
+  test_func_in_ns_with_same_name(q, ctxt);
+  test_function_with_class(q, ctxt);
+  test_fucntions_with_int_class_alias(q, ctxt);
+  return 0;
+}