Skip to content

Commit a784071

Browse files
alexeyvoronov-intelbader
alexeyvoronov-intel
authored and
bader
committed
[SYCL] Fix the integration header generation for kernel with typedefs.
Added printing policy with the canonical types flag enabled to generation declarations of templated kernel function types and specializations of KernelInfo for kernel function types. Before this change for single_task/parallel_for<class kernel<T>> where T is type alias (typedef/using) in the integration header remained the alias without defining what it is. Example: //// SYCL code using type_t = int; template<type_t N> class kernel; single_task<class kernel<0>>(... //// generated integration header // This is auto-generated SYCL integration header. // Forward declarations of templated kernel function types: template <type_t N> class kernel; Signed-off-by: Alexey Voronov <[email protected]>
1 parent f2ab989 commit a784071

File tree

5 files changed

+98
-14
lines changed

5 files changed

+98
-14
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

100755100644
Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1299,6 +1299,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
12991299
// print declaration into a string:
13001300
PrintingPolicy P(D->getASTContext().getLangOpts());
13011301
P.adjustForCPlusPlusFwdDecl();
1302+
P.PrintCanonicalTypes = true;
13021303
std::string S;
13031304
llvm::raw_string_ostream SO(S);
13041305
D->print(SO, P);
@@ -1511,8 +1512,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
15111512
O << "', '" << c;
15121513
O << "'> {\n";
15131514
} else {
1515+
LangOptions LO;
1516+
PrintingPolicy P(LO);
1517+
P.PrintCanonicalTypes = true;
15141518
O << "template <> struct KernelInfo<"
1515-
<< eraseAnonNamespace(K.NameType.getAsString()) << "> {\n";
1519+
<< eraseAnonNamespace(K.NameType.getAsString(P)) << "> {\n";
15161520
}
15171521
O << " DLL_LOCAL\n";
15181522
O << " static constexpr const char* getName() { return \"" << K.Name

‎clang/test/CodeGenSYCL/int_header1.cpp

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,15 @@
22
// RUN: FileCheck -input-file=%t.h %s
33

44
// CHECK:template <> struct KernelInfo<class KernelName> {
5-
// CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
6-
// CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
7-
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::nm2::KernelName0>> {
8-
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::KernelName1>> {
9-
// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::nm2::KernelName0>> {
10-
// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::KernelName1>> {
11-
// CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
12-
// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
13-
// CHECK:template <> struct KernelInfo<::nm1::KernelName8< ::nm1::nm2::C>> {
5+
// CHECK:template <> struct KernelInfo<class nm1::nm2::KernelName0> {
6+
// CHECK:template <> struct KernelInfo<class nm1::KernelName1> {
7+
// CHECK:template <> struct KernelInfo<class nm1::KernelName3<class nm1::nm2::KernelName0>> {
8+
// CHECK:template <> struct KernelInfo<class nm1::KernelName3<class nm1::KernelName1>> {
9+
// CHECK:template <> struct KernelInfo<class nm1::KernelName4<class nm1::nm2::KernelName0>> {
10+
// CHECK:template <> struct KernelInfo<class nm1::KernelName4<class nm1::KernelName1>> {
11+
// CHECK:template <> struct KernelInfo<class nm1::KernelName3<class KernelName5>> {
12+
// CHECK:template <> struct KernelInfo<class nm1::KernelName4<class KernelName7>> {
13+
// CHECK:template <> struct KernelInfo<class nm1::KernelName8<class nm1::nm2::C>> {
1414
// CHECK:template <> struct KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>> {
1515

1616
// This test checks if the SYCL device compiler is able to generate correct

‎clang/test/CodeGenSYCL/integration_header.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -48,9 +48,9 @@
4848
// CHECK-NEXT: };
4949
//
5050
// CHECK: template <> struct KernelInfo<class first_kernel> {
51-
// CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
52-
// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X> >> {
53-
// CHECK: template <> struct KernelInfo<::fourth_kernel< ::template_arg_ns::namespaced_arg<1> >> {
51+
// CHECK: template <> struct KernelInfo<class second_namespace::second_kernel<char>> {
52+
// CHECK: template <> struct KernelInfo<class third_kernel<1, int, struct point<struct X> >> {
53+
// CHECK: template <> struct KernelInfo<class fourth_kernel<struct template_arg_ns::namespaced_arg<1> >> {
5454

5555
#include "sycl.hpp"
5656

‎clang/test/CodeGenSYCL/kernel_functor.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -168,7 +168,7 @@ int main() {
168168
cl::sycl::detail::KernelInfo<Functor1>::getName();
169169
// CHECK: Functor1
170170
cl::sycl::detail::KernelInfo<ns::Functor2>::getName();
171-
// CHECK: ::ns::Functor2
171+
// CHECK: ns::Functor2
172172
cl::sycl::detail::KernelInfo<TmplFunctor<int>>::getName();
173173
// CHECK: TmplFunctor<int>
174174
cl::sycl::detail::KernelInfo<TmplConstFunctor<int>>::getName();

‎clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
1+
// RUN: %clang -I %S/Inputs --sycl -Xclang -fsycl-int-header=%t.h %s -c -o kernel.spv
2+
// RUN: FileCheck -input-file=%t.h %s
3+
4+
// CHECK: // Forward declarations of templated kernel function types:
5+
// CHECK-NEXT: template <typename T, typename T2, long N, unsigned long M> struct functor1;
6+
// CHECK-NEXT: template <typename T, typename T2, long N, unsigned long M> struct functor2;
7+
// CHECK-NEXT: template <typename T, typename T2> struct functor3;
8+
//
9+
// CHECK: // Specializations of KernelInfo for kernel function types:
10+
// CHECK: template <> struct KernelInfo<struct functor1<long, unsigned long, 0, 1>> {
11+
// CHECK: template <> struct KernelInfo<struct functor2<long, unsigned long, 0, 1>> {
12+
// CHECK: template <> struct KernelInfo<struct functor3<int, int>> {
13+
// CHECK: template <> struct KernelInfo<struct functor3<long, int>> {
14+
// CHECK: template <> struct KernelInfo<struct functor3<int, unsigned long>> {
15+
// CHECK: template <> struct KernelInfo<struct functor3<long, float>> {
16+
// CHECK: template <> struct KernelInfo<struct functor3<float, unsigned long>> {
17+
// CHECK: template <> struct KernelInfo<struct functor3<long, unsigned long>> {
18+
19+
#include "sycl.hpp"
20+
21+
template <typename KernelName, typename KernelType>
22+
__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
23+
kernelFunc();
24+
}
25+
26+
typedef signed long int signed_integer_t;
27+
28+
using unsigned_integer_t = unsigned long int;
29+
30+
template <typename T, typename T2, signed long int N, unsigned long int M>
31+
struct functor1 { void operator()() {} };
32+
33+
template <typename T, typename T2, signed_integer_t N, unsigned_integer_t M>
34+
struct functor2 { void operator()() {} };
35+
36+
template <typename T, typename T2>
37+
struct functor3 { void operator()() {} };
38+
39+
template <typename T>
40+
struct functor3<signed_integer_t, T> { void operator()() {} };
41+
42+
template <typename T>
43+
struct functor3<T, unsigned_integer_t> { void operator()() {} };
44+
45+
template <>
46+
struct functor3<signed_integer_t, float> { void operator()() {} };
47+
48+
template <>
49+
struct functor3<float, unsigned_integer_t> { void operator()() {} };
50+
51+
template <>
52+
struct functor3<signed_integer_t, unsigned_integer_t> { void operator()() {} };
53+
54+
int main() {
55+
functor1<signed long int, unsigned long int, 0L, 1UL> Functor1;
56+
kernel_single_task<decltype(Functor1)>(Functor1);
57+
58+
functor2<signed_integer_t, unsigned_integer_t, 0L, 1UL> Functor2;
59+
kernel_single_task<decltype(Functor2)>(Functor2);
60+
61+
functor3<int, int> Functor3;
62+
kernel_single_task<decltype(Functor3)>(Functor3);
63+
64+
functor3<signed_integer_t, int> Functor4;
65+
kernel_single_task<decltype(Functor4)>(Functor4);
66+
67+
functor3<int, unsigned_integer_t> Functor5;
68+
kernel_single_task<decltype(Functor5)>(Functor5);
69+
70+
functor3<signed_integer_t, float> Functor6;
71+
kernel_single_task<decltype(Functor6)>(Functor6);
72+
73+
functor3<float, unsigned_integer_t> Functor7;
74+
kernel_single_task<decltype(Functor7)>(Functor7);
75+
76+
functor3<signed_integer_t, unsigned_integer_t> Functor8;
77+
kernel_single_task<decltype(Functor8)>(Functor8);
78+
79+
return 0;
80+
}

0 commit comments

Comments
 (0)