Skip to content

Commit 5d22a8e

Browse files
alexeyvoronov-intelbader
alexeyvoronov-intel
authored and
bader
committed
[SYCL] Add the global namespace specifier for the kernel names if necessary.
Added a fix for the integration header generation - instead of policy PrintCanonicalTypes, the new policy SuppressTypedefs is used. Policy SuppressTypedefs only works for aliases and does not violate namespaces for non-builtin types. The problem caused if a user will use a type from namespace cl, sycl or detail. The compiler will try to found the type inside SYCL internal namespaces, not the user namespaces. E.g.: /* integration_header.hpp */ ... // forward declarations for kernel names namespace sycl { struct B; } ... // KernelInfo namespace cl { namespace sycl { detail { ... // before the patch, does not work template <> struct KernelInfo<struct sycl::B> { // after the patch, works template <> struct KernelInfo<::sycl::B> { ... }}} ... C++ looks for namespaces from current to global, and in this example the compiler finds the sycl namespace before it's necessary. Enforcing search start from the global namespace resolves this problem. Signed-off-by: Alexey Voronov <[email protected]>
1 parent 63e2b19 commit 5d22a8e

File tree

6 files changed

+128
-69
lines changed

6 files changed

+128
-69
lines changed

‎clang/include/clang/AST/PrettyPrinter.h

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ struct PrintingPolicy {
4242
SuppressScope(false), SuppressUnwrittenScope(false),
4343
SuppressInitializers(false), ConstantArraySizeAsWritten(false),
4444
AnonymousTagLocations(true), SuppressStrongLifetime(false),
45-
SuppressLifetimeQualifiers(false),
45+
SuppressLifetimeQualifiers(false), SuppressTypedefs(false),
4646
SuppressTemplateArgsInCXXConstructors(false), Bool(LO.Bool),
4747
Restrict(LO.C99), Alignof(LO.CPlusPlus11), UnderscoreAlignof(LO.C11),
4848
UseVoidForZeroParams(!LO.CPlusPlus), TerseOutput(false),
@@ -159,6 +159,17 @@ struct PrintingPolicy {
159159
/// When true, suppress printing of lifetime qualifier in ARC.
160160
unsigned SuppressLifetimeQualifiers : 1;
161161

162+
/// When true prints a canonical type instead of an alias. E.g.
163+
/// \code
164+
/// using SizeT = int;
165+
/// template<SizeT N> class C;
166+
/// \endcode
167+
/// will be printed as
168+
/// \code
169+
/// template<int N> class C;
170+
/// \endcode
171+
unsigned SuppressTypedefs : 1;
172+
162173
/// When true, suppresses printing template arguments in names of C++
163174
/// constructors.
164175
unsigned SuppressTemplateArgsInCXXConstructors : 1;

‎clang/lib/AST/TypePrinter.cpp

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -292,6 +292,14 @@ void TypePrinter::printBefore(const Type *T,Qualifiers Quals, raw_ostream &OS) {
292292
if (Policy.SuppressSpecifiers && T->isSpecifierType())
293293
return;
294294

295+
if (Policy.SuppressTypedefs && (T->getTypeClass() == Type::Typedef)) {
296+
QualType UnderlyingType = T->getCanonicalTypeInternal();
297+
SplitQualType Split = splitAccordingToPolicy(UnderlyingType, Policy);
298+
Qualifiers FullQuals = Quals + Split.Quals;
299+
printBefore(Split.Ty, FullQuals, OS);
300+
return;
301+
}
302+
295303
SaveAndRestore<bool> PrevPHIsEmpty(HasEmptyPlaceHolder);
296304

297305
// Print qualifiers as appropriate.
@@ -942,7 +950,6 @@ void TypePrinter::printFunctionNoProtoAfter(const FunctionNoProtoType *T,
942950
}
943951

944952
void TypePrinter::printTypeSpec(NamedDecl *D, raw_ostream &OS) {
945-
946953
// Compute the full nested-name-specifier for this type.
947954
// In C, this will always be empty except when the type
948955
// being printed is anonymous within other Record.
@@ -1323,7 +1330,8 @@ void TypePrinter::printElaboratedBefore(const ElaboratedType *T,
13231330
if (T->getKeyword() != ETK_None)
13241331
OS << " ";
13251332
NestedNameSpecifier *Qualifier = T->getQualifier();
1326-
if (Qualifier)
1333+
if (Qualifier && !(Policy.SuppressTypedefs &&
1334+
T->getNamedType()->getTypeClass() == Type::Typedef))
13271335
Qualifier->print(OS, Policy);
13281336
}
13291337

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1423,7 +1423,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
14231423
// print declaration into a string:
14241424
PrintingPolicy P(D->getASTContext().getLangOpts());
14251425
P.adjustForCPlusPlusFwdDecl();
1426-
P.PrintCanonicalTypes = true;
1426+
P.SuppressTypedefs = true;
14271427
std::string S;
14281428
llvm::raw_string_ostream SO(S);
14291429
D->print(SO, P);
@@ -1638,7 +1638,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
16381638
} else {
16391639
LangOptions LO;
16401640
PrintingPolicy P(LO);
1641-
P.PrintCanonicalTypes = true;
1641+
P.SuppressTypedefs = true;
16421642
O << "template <> struct KernelInfo<"
16431643
<< eraseAnonNamespace(K.NameType.getAsString(P)) << "> {\n";
16441644
}

‎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<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>> {
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>> {
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<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> >> {
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> >> {
5454

5555
#include "sycl.hpp"
5656

‎clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp

Lines changed: 92 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -1,80 +1,120 @@
11
// RUN: %clang -I %S/Inputs --sycl -Xclang -fsycl-int-header=%t.h %s -c -o kernel.spv
22
// RUN: FileCheck -input-file=%t.h %s
33

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-
194
#include "sycl.hpp"
205

216
template <typename KernelName, typename KernelType>
22-
__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
7+
__attribute__((sycl_kernel)) void single_task(KernelType kernelFunc) {
238
kernelFunc();
249
}
2510

26-
typedef signed long int signed_integer_t;
27-
28-
using unsigned_integer_t = unsigned long int;
11+
struct dummy_functor {
12+
void operator()() {}
13+
};
2914

30-
template <typename T, typename T2, signed long int N, unsigned long int M>
31-
struct functor1 { void operator()() {} };
15+
typedef int int_t;
16+
using uint_t = unsigned int;
3217

33-
template <typename T, typename T2, signed_integer_t N, unsigned_integer_t M>
34-
struct functor2 { void operator()() {} };
18+
typedef const int cint_t;
19+
using cuint_t = const unsigned int;
3520

36-
template <typename T, typename T2>
37-
struct functor3 { void operator()() {} };
21+
namespace space {
22+
typedef long long_t;
23+
using ulong_t = unsigned long;
24+
typedef const long clong_t;
25+
using culong_t = const unsigned long;
26+
} // namespace space
3827

39-
template <typename T>
40-
struct functor3<signed_integer_t, T> { void operator()() {} };
28+
// non-type template arguments cases
29+
// CHECK: template <int N, unsigned int M> struct kernel_name1;
30+
template <int N, unsigned int M>
31+
struct kernel_name1 {};
4132

42-
template <typename T>
43-
struct functor3<T, unsigned_integer_t> { void operator()() {} };
33+
// CHECK: template <int N, unsigned int M> struct kernel_name1v1;
34+
template <int_t N, uint_t M>
35+
struct kernel_name1v1 {};
4436

45-
template <>
46-
struct functor3<signed_integer_t, float> { void operator()() {} };
37+
// CHECK: template <long N, unsigned long M> struct kernel_name1v2;
38+
template <space::long_t N, space::ulong_t M>
39+
struct kernel_name1v2 {};
4740

48-
template <>
49-
struct functor3<float, unsigned_integer_t> { void operator()() {} };
41+
// CHECK: template <typename T, typename T2> struct kernel_name2;
42+
template <typename T, typename T2>
43+
struct kernel_name2;
44+
45+
// CHECK: class A;
46+
class A {};
47+
namespace space {
48+
// CHECK: namespace space {
49+
// CHECK-NEXT: class B;
50+
// CHECK-NEXT: }
51+
class B {};
52+
using a_t = A;
53+
using b_t = B;
54+
} // namespace space
55+
56+
// partial template specialization cases
57+
template <typename T>
58+
struct kernel_name2<int_t, T> {};
5059

51-
template <>
52-
struct functor3<signed_integer_t, unsigned_integer_t> { void operator()() {} };
60+
template <typename T>
61+
struct kernel_name2<cint_t, T> {};
5362

54-
int main() {
55-
functor1<signed long int, unsigned long int, 0L, 1UL> Functor1;
56-
kernel_single_task<decltype(Functor1)>(Functor1);
63+
template <typename T>
64+
struct kernel_name2<space::long_t, T> {};
5765

58-
functor2<signed_integer_t, unsigned_integer_t, 0L, 1UL> Functor2;
59-
kernel_single_task<decltype(Functor2)>(Functor2);
66+
template <typename T>
67+
struct kernel_name2<const space::long_t, T> {};
6068

61-
functor3<int, int> Functor3;
62-
kernel_single_task<decltype(Functor3)>(Functor3);
69+
template <typename T>
70+
struct kernel_name2<volatile space::clong_t, T> {};
6371

64-
functor3<signed_integer_t, int> Functor4;
65-
kernel_single_task<decltype(Functor4)>(Functor4);
72+
template <typename T>
73+
struct kernel_name2<space::a_t, T> {};
6674

67-
functor3<int, unsigned_integer_t> Functor5;
68-
kernel_single_task<decltype(Functor5)>(Functor5);
75+
template <typename T>
76+
struct kernel_name2<space::b_t, T> {};
6977

70-
functor3<signed_integer_t, float> Functor6;
71-
kernel_single_task<decltype(Functor6)>(Functor6);
78+
// full template specialization cases
79+
template <>
80+
struct kernel_name2<int_t, const uint_t> {};
7281

73-
functor3<float, unsigned_integer_t> Functor7;
74-
kernel_single_task<decltype(Functor7)>(Functor7);
82+
template <>
83+
struct kernel_name2<space::clong_t, volatile space::culong_t> {};
7584

76-
functor3<signed_integer_t, unsigned_integer_t> Functor8;
77-
kernel_single_task<decltype(Functor8)>(Functor8);
85+
template <>
86+
struct kernel_name2<space::a_t, volatile space::b_t> {};
7887

88+
int main() {
89+
dummy_functor f;
90+
// non-type template arguments
91+
// CHECK: template <> struct KernelInfo<::kernel_name1<1, 1>> {
92+
single_task<kernel_name1<1, 1>>(f);
93+
// CHECK: template <> struct KernelInfo<::kernel_name1v1<1, 1>> {
94+
single_task<kernel_name1v1<1, 1>>(f);
95+
// CHECK: template <> struct KernelInfo<::kernel_name1v2<1, 1>> {
96+
single_task<kernel_name1v2<1, 1>>(f);
97+
// partial template specialization
98+
// CHECK: template <> struct KernelInfo<::kernel_name2<int, int>> {
99+
single_task<kernel_name2<int_t, int>>(f);
100+
// CHECK: template <> struct KernelInfo<::kernel_name2<const int, char>> {
101+
single_task<kernel_name2<cint_t, char>>(f);
102+
// CHECK: template <> struct KernelInfo<::kernel_name2<long, float>> {
103+
single_task<kernel_name2<space::long_t, float>>(f);
104+
// CHECK: template <> struct KernelInfo<::kernel_name2<const long, ::A>> {
105+
single_task<kernel_name2<const space::long_t, space::a_t>>(f);
106+
// CHECK: template <> struct KernelInfo<::kernel_name2<const volatile long, const ::space::B>> {
107+
single_task<kernel_name2<volatile space::clong_t, const space::b_t>>(f);
108+
// CHECK: template <> struct KernelInfo<::kernel_name2< ::A, long>> {
109+
single_task<kernel_name2<space::a_t, space::long_t>>(f);
110+
// CHECK: template <> struct KernelInfo<::kernel_name2< ::space::B, int>> {
111+
single_task<kernel_name2<space::b_t, int_t>>(f);
112+
// full template specialization
113+
// CHECK: template <> struct KernelInfo<::kernel_name2<int, const unsigned int>> {
114+
single_task<kernel_name2<int_t, const uint_t>>(f);
115+
// CHECK: template <> struct KernelInfo<::kernel_name2<const long, volatile const unsigned long>> {
116+
single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
117+
// CHECK: template <> struct KernelInfo<::kernel_name2< ::A, volatile ::space::B>> {
118+
single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
79119
return 0;
80120
}

0 commit comments

Comments
 (0)