[SYCL] Enable use of variadic templates in SYCL kernel names.

clang/lib/AST/QualTypeNames.cpp

@@ -109,6 +109,11 @@ static bool getFullyQualifiedTemplateArgument(const ASTContext &Ctx,
       Arg = TemplateArgument(QTFQ);
       Changed = true;
     }
+  } else if (Arg.getKind() == TemplateArgument::Pack) {
+    for (const auto &It : Arg.pack_elements()) {
+      Changed |= getFullyQualifiedTemplateArgument(
+          Ctx, const_cast<TemplateArgument &>(It), WithGlobalNsPrefix);
+    }
   }
   return Changed;
 }

clang/test/CodeGenSYCL/int_header1.cpp

@@ -11,6 +11,7 @@
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::KernelName1>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName8< ::nm1::nm2::C>> {
 // CHECK:template <> struct KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>> {
 
 // This test checks if the SYCL device compiler is able to generate correct
@@ -34,6 +35,7 @@ namespace nm1 {
 
   template <typename T> class KernelName3;
   template <typename T> class KernelName4;
+  template <typename... T> class KernelName8;
 
   template <> class KernelName3<nm1::nm2::KernelName0>;
   template <> class KernelName3<KernelName1>;
@@ -142,6 +144,12 @@ struct MyWrapper {
     kernel_single_task<nm1::KernelName4<class KernelName7>>(
       [=]() { acc.use(); });
 
+    // TPITD
+    // a defined template pack specialization class with defined class
+    // as argument declared in a namespace at translation unit scope
+    kernel_single_task<nm1::KernelName8<nm1::nm2::C>>(
+      [=]() { acc.use(); });
+
 #ifdef TDLI
     // TODO unexpected compilation error when host code + integration header
     // is compiled TDLI a defined template specialization class with
@@ -202,6 +210,7 @@ int main() {
   KernelInfo<class nm1::KernelName4<class nm1::KernelName1>>::getName();
   KernelInfo<class nm1::KernelName3<class KernelName5>>::getName();
   KernelInfo<class nm1::KernelName4<class KernelName7>>::getName();
+  KernelInfo<class nm1::KernelName8<nm1::nm2::C>>::getName();
   KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>>::getName();
 #endif //__SYCL_DEVICE_ONLY__
 }

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
@@ -36,12 +41,17 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 8 },
 // 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, 6112, 4 },
+// 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<::fourth_kernel< ::template_arg_ns::namespaced_arg<1> >> {
 
 #include "sycl.hpp"
 
@@ -60,6 +70,14 @@ class second_kernel;
 template <int a, typename T1, typename T2>
 class third_kernel;
 
+namespace template_arg_ns {
+template <int DimX>
+struct namespaced_arg {};
+} // namespace template_arg_ns
+
+template <typename... Ts>
+class fourth_kernel;
+
 int main() {
 
   cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc1;
@@ -98,6 +116,11 @@ int main() {
     }
   });
 
+    kernel_single_task<class fourth_kernel<template_arg_ns::namespaced_arg<1>>>([=]() {
+      if (i == 13) {
+        acc2.use();
+      }
+  });
+
   return 0;
 }
-

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,27 @@ struct Wrapper {
       });
       ++NumTestCases;
 #endif
+      // TPITD
+      // an incomplete vatiadic template specialization class in a namespace at
+      // translation unit scope with a defined class as argument declared in
+      // a namespace at translation unit scope
+      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;
+
+      // TPITD
+      // an incomplete vatiadic template specialization class in the global
+      // namespace at translation unit scope with a defined class as argument
+      // declared in a namespace at translation unit scope
+      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];
   }
@@ -246,4 +273,3 @@ int main() {
   std::cout << (pass ? "pass" : "FAIL") << "\n";
   return pass ? 0 : 1;
 }
-