diff --git a/clang/lib/AST/QualTypeNames.cpp b/clang/lib/AST/QualTypeNames.cpp index ce01f66297bce..ff5684c023309 100644 --- a/clang/lib/AST/QualTypeNames.cpp +++ b/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; } diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 24c5d13ce8d99..d2bf8f16848f7 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/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__ } diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index f911b4cd9941f..f15b6df67711e 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 @@ -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; } - diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp index 4e4d2e835ef55..f09b8d241a135 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,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; } -