diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c7301d5970b6e..789da27fc0091 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -736,6 +736,20 @@ static std::string eraseAnonNamespace(std::string S) {
   return S;
 }
 
+// Removes all "class" and "struct" substrings from given string
+static std::string eraseClassAndStruct(std::string S) {
+  const char S1[] = "class ";
+  const char S2[] = "struct ";
+
+  for (auto Pos = S.find(S1); Pos != StringRef::npos; Pos = S.find(S1, Pos))
+    S.erase(Pos, sizeof(S1) - 1);
+
+  for (auto Pos = S.find(S2); Pos != StringRef::npos; Pos = S.find(S2, Pos))
+    S.erase(Pos, sizeof(S2) - 1);
+
+  return S;
+}
+
 // Creates a mangled kernel name for given kernel name type
 static std::string constructKernelName(QualType KernelNameType,
                                        ASTContext &AC) {
@@ -761,8 +775,7 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) {
   assert(TemplateArgs && "No template argument info");
   // The first template argument always describes the kernel name - whether
   // it is lambda or functor.
-  QualType KernelNameType = TypeName::getFullyQualifiedType(
-      TemplateArgs->get(0).getAsType(), getASTContext(), true);
+  QualType KernelNameType = TemplateArgs->get(0).getAsType();
   std::string Name = constructKernelName(KernelNameType, getASTContext());
   populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE);
   FunctionDecl *SYCLKernel =
@@ -1025,7 +1038,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   for (const KernelDesc &K : KernelDescs) {
     const size_t N = K.Params.size();
     O << "template <> struct KernelInfo<"
-      << eraseAnonNamespace(K.NameType.getAsString()) << "> {\n";
+      << eraseClassAndStruct(eraseAnonNamespace(K.NameType.getAsString()))
+      << "> {\n";
     O << "  static constexpr const char* getName() { return \"" << K.Name
       << "\"; }\n";
     O << "  static constexpr unsigned getNumParams() { return " << N << "; }\n";
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index 9e457819562d9..cf46052437a46 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -10,12 +10,17 @@
 // CHECK-NEXT: struct X;
 // CHECK-NEXT: template <typename T> struct point;
 // CHECK-NEXT: template <int a, typename T1, typename T2> class third_kernel;
+// CHECK-NEXT: namespace template_arg_ns {
+// CHECK-NEXT: template <int DimX> struct namespaced_arg;
+// CHECK-NEXT: }
+// CHECK-NEXT: template <typename ...Ts> class fourth_kernel;
 //
 // CHECK: static constexpr
 // CHECK-NEXT: const char* const kernel_names[] = {
 // CHECK-NEXT:   "_ZTSZ4mainE12first_kernel",
 // CHECK-NEXT:   "_ZTSN16second_namespace13second_kernelIcEE",
 // CHECK-NEXT:   "_ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE"
+// CHECK-NEXT:   "_ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE"
 // CHECK-NEXT: };
 //
 // CHECK: static constexpr
@@ -41,12 +46,19 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 5 },
 // CHECK-EMPTY:
+// CHECK-NEXT:   //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 2016, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 5 },
+// CHECK-EMPTY:
 // CHECK-NEXT: };
 //
 // CHECK: template <class KernelNameType> struct KernelInfo;
-// CHECK: template <> struct KernelInfo<class first_kernel> {
-// CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
-// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X> >> {
+// CHECK: template <> struct KernelInfo<first_kernel> {
+// CHECK: template <> struct KernelInfo<second_namespace::second_kernel<char>> {
+// CHECK: template <> struct KernelInfo<third_kernel<1, int, point<X> >> {
+// CHECK: template <> struct KernelInfo<fourth_kernel<template_arg_ns::namespaced_arg<1> >> {
 
 namespace cl {
 namespace sycl {
@@ -124,6 +136,14 @@ class second_kernel;
 template <int a, typename T1, typename T2>
 class third_kernel;
 
+namespace template_arg_ns {
+  template <int DimX>
+  struct namespaced_arg {};
+}
+
+template <typename ...Ts>
+class fourth_kernel;
+
 int main() {
 
   cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc1;
@@ -157,6 +177,11 @@ int main() {
       acc2.use();
     }
   });
+  kernel_single_task<class fourth_kernel<template_arg_ns::namespaced_arg<1>>>([=]() {
+    if (i == 13) {
+      acc2.use();
+    }
+  });
 
   return 0;
 }
diff --git a/clang/test/CodeGenSYCL/kernel_functor.cpp b/clang/test/CodeGenSYCL/kernel_functor.cpp
index 56b9b34c2dbeb..c37f6a343b1de 100644
--- a/clang/test/CodeGenSYCL/kernel_functor.cpp
+++ b/clang/test/CodeGenSYCL/kernel_functor.cpp
@@ -168,7 +168,7 @@ int main() {
   cl::sycl::detail::KernelInfo<Functor1>::getName();
   // CHECK: Functor1
   cl::sycl::detail::KernelInfo<ns::Functor2>::getName();
-  // CHECK: ::ns::Functor2
+  // CHECK: ns::Functor2
   cl::sycl::detail::KernelInfo<TmplFunctor<int>>::getName();
   // CHECK: TmplFunctor<int>
   cl::sycl::detail::KernelInfo<TmplConstFunctor<int>>::getName();
diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp
index 4e4d2e835ef55..9de14e1ba99e6 100644
--- a/sycl/test/regression/kernel_name_class.cpp
+++ b/sycl/test/regression/kernel_name_class.cpp
@@ -21,6 +21,8 @@ namespace nm1 {
 namespace nm2 {
 class C {};
 class KernelName0 : public C {};
+
+template <int X> class KernelName11 {};
 } // namespace nm2
 
 class KernelName1;
@@ -34,8 +36,12 @@ template <> class KernelName3<KernelName1>;
 template <> class KernelName4<nm1::nm2::KernelName0> {};
 template <> class KernelName4<KernelName1> {};
 
+template <typename... Ts> class KernelName10;
+
 } // namespace nm1
 
+template <typename... Ts> class KernelName12;
+
 static int NumTestCases = 0;
 
 namespace nm3 {
@@ -233,6 +239,19 @@ struct Wrapper {
       });
       ++NumTestCases;
 #endif
+
+      deviceQueue.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task<nm1::KernelName10<nm1::nm2::KernelName11<10>>>([=]() { acc[0] += GOLD; });
+      });
+      ++NumTestCases;
+
+      deviceQueue.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task<KernelName12<nm1::nm2::KernelName11<10>>>([=]() { acc[0] += GOLD; });
+      });
+      ++NumTestCases;
+
     }
     return arr[0];
   }