Skip to content

Commit 964bb7b

Browse files
elizabethandrewselizabethandrews
committed
Add support to handle enums when KernelNameType is templated
using a type which is in turn templated using enum. Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
1 parent 08f8656 commit 964bb7b

File tree

3 files changed

+54
-17
lines changed

3 files changed

+54
-17
lines changed

‎clang/include/clang/Sema/Sema.h

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -319,7 +319,8 @@ class SYCLIntegrationHeader {
319319
};
320320

321321
public:
322-
SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport);
322+
SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport,
323+
Sema &S);
323324

324325
/// Emits contents of the header into given stream.
325326
void emit(raw_ostream &Out);
@@ -424,6 +425,8 @@ class SYCLIntegrationHeader {
424425

425426
/// Whether header is generated with unnamed lambda support
426427
bool UnnamedLambdaSupport;
428+
429+
Sema &S;
427430
};
428431

429432
/// Keeps track of expected type during expression parsing. The type is tied to
@@ -12584,7 +12587,7 @@ class Sema final {
1258412587
SYCLIntegrationHeader &getSyclIntegrationHeader() {
1258512588
if (SyclIntHeader == nullptr)
1258612589
SyclIntHeader = std::make_unique<SYCLIntegrationHeader>(
12587-
getDiagnostics(), getLangOpts().SYCLUnnamedLambda);
12590+
getDiagnostics(), getLangOpts().SYCLUnnamedLambda, *this);
1258812591
return *SyclIntHeader.get();
1258912592
}
1259012593

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 19 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1863,6 +1863,9 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
18631863
ArrayRef<TemplateArgument> Args,
18641864
const PrintingPolicy &P);
18651865

1866+
static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
1867+
const PrintingPolicy &TypePolicy);
1868+
18661869
static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
18671870
TemplateArgument Arg, const PrintingPolicy &P) {
18681871
switch (Arg.getKind()) {
@@ -1888,8 +1891,7 @@ static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
18881891
TypePolicy.SuppressTypedefs = true;
18891892
TypePolicy.SuppressTagKeyword = true;
18901893
QualType T = Arg.getAsType();
1891-
QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);
1892-
ArgOS << FullyQualifiedType.getAsString(TypePolicy);
1894+
ArgOS << getKernelNameTypeString(T, Ctx, TypePolicy);
18931895
break;
18941896
}
18951897
default:
@@ -1918,36 +1920,36 @@ static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
19181920
ArgOS << ">";
19191921
}
19201922

1921-
static std::string getKernelNameTypeString(QualType T) {
1923+
static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
1924+
const PrintingPolicy &TypePolicy) {
1925+
1926+
QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);
19221927

19231928
const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
19241929

19251930
if (!RD)
1926-
return getCPPTypeString(T);
1931+
return eraseAnonNamespace(FullyQualifiedType.getAsString(TypePolicy));
19271932

19281933
// If kernel name type is a template specialization with enum type
19291934
// template parameters, enumerators in name type string should be
19301935
// replaced with their underlying value since the enum definition
19311936
// is not visible in integration header.
19321937
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
1933-
LangOptions LO;
1934-
PrintingPolicy P(LO);
1935-
P.SuppressTypedefs = true;
19361938
SmallString<64> Buf;
19371939
llvm::raw_svector_ostream ArgOS(Buf);
19381940

19391941
// Print template class name
1940-
TSD->printQualifiedName(ArgOS, P, /*WithGlobalNsPrefix*/ true);
1942+
TSD->printQualifiedName(ArgOS, TypePolicy, /*WithGlobalNsPrefix*/ true);
19411943

19421944
// Print template arguments substituting enumerators
19431945
ASTContext &Ctx = RD->getASTContext();
19441946
const TemplateArgumentList &Args = TSD->getTemplateArgs();
1945-
printTemplateArguments(Ctx, ArgOS, Args.asArray(), P);
1947+
printTemplateArguments(Ctx, ArgOS, Args.asArray(), TypePolicy);
19461948

19471949
return eraseAnonNamespace(ArgOS.str().str());
19481950
}
19491951

1950-
return getCPPTypeString(T);
1952+
return eraseAnonNamespace(FullyQualifiedType.getAsString(TypePolicy));
19511953
}
19521954

19531955
void SYCLIntegrationHeader::emit(raw_ostream &O) {
@@ -2066,9 +2068,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
20662068
O << "', '" << c;
20672069
O << "'> {\n";
20682070
} else {
2069-
2071+
LangOptions LO;
2072+
PrintingPolicy P(LO);
2073+
P.SuppressTypedefs = true;
20702074
O << "template <> struct KernelInfo<"
2071-
<< getKernelNameTypeString(K.NameType) << "> {\n";
2075+
<< getKernelNameTypeString(K.NameType, S.getASTContext(), P) << "> {\n";
20722076
}
20732077
O << " DLL_LOCAL\n";
20742078
O << " static constexpr const char* getName() { return \"" << K.Name
@@ -2137,8 +2141,9 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
21372141
}
21382142

21392143
SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
2140-
bool _UnnamedLambdaSupport)
2141-
: Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport) {}
2144+
bool _UnnamedLambdaSupport,
2145+
Sema &_S)
2146+
: Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport), S(_S) {}
21422147

21432148
// -----------------------------------------------------------------------------
21442149
// Utility class methods

‎clang/test/CodeGenSYCL/kernelname-enum.cpp

Lines changed: 30 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -79,6 +79,21 @@ class dummy_functor_7 {
7979
void operator()() {}
8080
};
8181

82+
namespace type_argument_template_enum {
83+
enum class E : int {
84+
A,
85+
B,
86+
C
87+
};
88+
}
89+
90+
template <typename T>
91+
class T1 {};
92+
template <type_argument_template_enum::E EnumValue>
93+
class T2 {};
94+
template <typename EnumType>
95+
class T3 {};
96+
8297
int main() {
8398

8499
dummy_functor_1<no_namespace_int::val_1> f1;
@@ -124,6 +139,14 @@ int main() {
124139
cgh.single_task(f8);
125140
});
126141

142+
q.submit([&](cl::sycl::handler &cgh) {
143+
cgh.single_task<T1<T2<type_argument_template_enum::E::A>>>([=]() {});
144+
});
145+
146+
q.submit([&](cl::sycl::handler &cgh) {
147+
cgh.single_task<T1<T3<type_argument_template_enum::E>>>([=]() {});
148+
});
149+
127150
return 0;
128151
}
129152

@@ -145,7 +168,11 @@ int main() {
145168
// CHECK: enum unscoped_enum : int;
146169
// CHECK: template <unscoped_enum EnumType> class dummy_functor_6;
147170
// CHECK: template <typename EnumType> class dummy_functor_7;
148-
171+
// CHECK: namespace type_argument_template_enum {
172+
// CHECK-NEXT: enum class E : int;
173+
// CHECK-NEXT: }
174+
// CHECK: template <type_argument_template_enum::E EnumValue> class T2;
175+
// CHECK: template <typename T> class T1;
149176
// CHECK: Specializations of KernelInfo for kernel function types:
150177
// CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
151178
// CHECK: template <> struct KernelInfo<::dummy_functor_2<(no_namespace_short)1>>
@@ -155,3 +182,5 @@ int main() {
155182
// CHECK: template <> struct KernelInfo<::dummy_functor_6<(unscoped_enum)0>>
156183
// CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>>
157184
// CHECK: template <> struct KernelInfo<::dummy_functor_7<::internal::namespace_short>>
185+
// CHECK: template <> struct KernelInfo<::T1<::T2<(type_argument_template_enum::E)0>>>
186+
// CHECK: template <> struct KernelInfo<::T1<::T3<::type_argument_template_enum::E>>>

0 commit comments

Comments
 (0)