-
Notifications
You must be signed in to change notification settings - Fork 802
[SYCL] Add support for kernel name types templated using enums. #1675
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from 4 commits
9715988
b337e43
c182c48
f6b9b59
07135ab
df0994a
46ef79d
6262d69
b59e1a3
0a83e16
52add49
9d07023
eead667
1064c45
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -1690,10 +1690,17 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D, | |
| PrintingPolicy P(D->getASTContext().getLangOpts()); | ||
| P.adjustForCPlusPlusFwdDecl(); | ||
| P.SuppressTypedefs = true; | ||
| P.SuppressUnwrittenScope = true; | ||
| std::string S; | ||
| llvm::raw_string_ostream SO(S); | ||
| D->print(SO, P); | ||
| O << SO.str() << ";\n"; | ||
| O << SO.str(); | ||
|
|
||
| if (const auto *ED = dyn_cast<EnumDecl>(D)) { | ||
| QualType T = ED->getIntegerType(); | ||
elizabethandrews marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| O << " : " << T.getAsString() << ";\n"; | ||
| } else | ||
| O << ";\n"; | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
|
|
||
| // print closing braces for namespaces if needed | ||
| for (unsigned I = 0; I < NamespaceCnt; ++I) | ||
|
|
@@ -1796,6 +1803,15 @@ void SYCLIntegrationHeader::emitForwardClassDecls( | |
| } | ||
| break; | ||
| } | ||
| case TemplateArgument::ArgKind::Integral: { | ||
| // Handle Kernel Name Type templated using enum. | ||
| QualType T = Arg.getIntegralType(); | ||
| if (const EnumType *ET = T->getAs<EnumType>()) { | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
| const EnumDecl *ED = ET->getDecl(); | ||
| emitFwdDecl(O, ED, KernelLocation); | ||
| } | ||
| break; | ||
| } | ||
| default: | ||
| break; // nop | ||
| } | ||
|
|
@@ -1822,6 +1838,90 @@ static std::string getCPPTypeString(QualType Ty) { | |
| return eraseAnonNamespace(Ty.getAsString(P)); | ||
| } | ||
|
|
||
| static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS, | ||
| ArrayRef<TemplateArgument> Args, | ||
| const PrintingPolicy &P, bool ParameterPack) { | ||
|
|
||
| if (!ParameterPack) | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
| ArgOS << "<"; | ||
|
|
||
| bool FirstArg = true; | ||
| const char *Comma = ", "; | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
|
|
||
| for (unsigned I = 0; I < Args.size(); I++) { | ||
| const TemplateArgument &Arg = Args[I]; | ||
| if (Arg.getKind() == TemplateArgument::ArgKind::Pack) { | ||
|
||
| if (Arg.pack_size() && !FirstArg) | ||
| ArgOS << Comma; | ||
| printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P, true); | ||
| } else if (Arg.getKind() == TemplateArgument::ArgKind::Integral) { | ||
| if (!FirstArg) | ||
| ArgOS << Comma; | ||
|
|
||
| QualType T = Arg.getIntegralType(); | ||
| const EnumType *ET = T->getAs<EnumType>(); | ||
|
|
||
| if (ET) { | ||
| const llvm::APSInt &Val = Arg.getAsIntegral(); | ||
| ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val; | ||
| } else { | ||
| Arg.print(P, ArgOS); | ||
| } | ||
| } else if (Arg.getKind() == TemplateArgument::ArgKind::Type) { | ||
| if (!FirstArg) | ||
| ArgOS << Comma; | ||
| LangOptions LO; | ||
| PrintingPolicy TypePolicy(LO); | ||
| TypePolicy.SuppressTypedefs = true; | ||
| TypePolicy.SuppressTagKeyword = true; | ||
| QualType T = Arg.getAsType(); | ||
| QualType FullyQualifiedType = | ||
| TypeName::getFullyQualifiedType(T, Ctx, true); | ||
| ArgOS << FullyQualifiedType.getAsString(TypePolicy); | ||
| } else { | ||
| if (!FirstArg) | ||
| ArgOS << Comma; | ||
| Arg.print(P, ArgOS); | ||
| } | ||
| FirstArg = false; | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
| } | ||
|
|
||
| if (!ParameterPack) | ||
| ArgOS << ">"; | ||
| } | ||
|
|
||
| static std::string getKernelNameTypeString(QualType T) { | ||
|
|
||
| const CXXRecordDecl *RD = T->getAsCXXRecordDecl(); | ||
|
|
||
| if (!RD) | ||
| return getCPPTypeString(T); | ||
|
|
||
| // If kernel name type is a template specialization with enum type | ||
| // template parameters, enumerators in name type string should be | ||
| // replaced with their underlying value since the enum definition | ||
| // is not visible in integration header. | ||
| if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) { | ||
| LangOptions LO; | ||
| PrintingPolicy P(LO); | ||
| P.SuppressTypedefs = true; | ||
| SmallString<64> Buf; | ||
| llvm::raw_svector_ostream ArgOS(Buf); | ||
|
|
||
| // Print template class name | ||
| TSD->printQualifiedName(ArgOS, P, true); | ||
|
|
||
| // Print template arguments substituting enumerators | ||
| ASTContext &Ctx = RD->getASTContext(); | ||
| const TemplateArgumentList &Args = TSD->getTemplateArgs(); | ||
| printArguments(Ctx, ArgOS, Args.asArray(), P, false); | ||
elizabethandrews marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
|
|
||
| return eraseAnonNamespace(ArgOS.str().str()); | ||
| } | ||
|
|
||
| return getCPPTypeString(T); | ||
| } | ||
|
|
||
| void SYCLIntegrationHeader::emit(raw_ostream &O) { | ||
| O << "// This is auto-generated SYCL integration header.\n"; | ||
| O << "\n"; | ||
|
|
@@ -1938,8 +2038,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { | |
| O << "', '" << c; | ||
| O << "'> {\n"; | ||
| } else { | ||
| O << "template <> struct KernelInfo<" << getCPPTypeString(K.NameType) | ||
| << "> {\n"; | ||
|
|
||
| O << "template <> struct KernelInfo<" | ||
| << getKernelNameTypeString(K.NameType) << "> {\n"; | ||
| } | ||
| O << " DLL_LOCAL\n"; | ||
| O << " static constexpr const char* getName() { return \"" << K.Name | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,119 @@ | ||
| // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only | ||
| // RUN: FileCheck -input-file=%t.h %s | ||
|
|
||
| #include "sycl.hpp" | ||
|
|
||
| enum class no_namespace_int : int { | ||
elizabethandrews marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| val_1, | ||
| val_2 | ||
| }; | ||
|
|
||
| enum class no_namespace_short : short { | ||
| val_1, | ||
| val_2 | ||
| }; | ||
|
|
||
| namespace internal { | ||
| enum class namespace_short : short { | ||
| val_1, | ||
| val_2 | ||
| }; | ||
| } | ||
|
|
||
| namespace { | ||
| enum class enum_in_anonNS : short { | ||
| val_1, | ||
| val_2 | ||
| }; | ||
| } | ||
|
|
||
| enum class no_type_set { | ||
| val_1, | ||
| val_2 | ||
| }; | ||
|
|
||
| template <no_namespace_int EnumType> | ||
| class dummy_functor_1 { | ||
| public: | ||
| void operator()() {} | ||
| }; | ||
|
|
||
| template <no_namespace_short EnumType> | ||
| class dummy_functor_2 { | ||
| public: | ||
| void operator()() {} | ||
| }; | ||
|
|
||
| template <internal::namespace_short EnumType> | ||
| class dummy_functor_3 { | ||
| public: | ||
| void operator()() {} | ||
| }; | ||
|
|
||
| template <enum_in_anonNS EnumType> | ||
| class dummy_functor_4 { | ||
| public: | ||
| void operator()() {} | ||
| }; | ||
|
|
||
| template <no_type_set EnumType> | ||
| class dummy_functor_5 { | ||
| public: | ||
| void operator()() {} | ||
| }; | ||
|
|
||
| int main() { | ||
|
|
||
| dummy_functor_1<no_namespace_int::val_1> f1; | ||
| dummy_functor_2<no_namespace_short::val_2> f2; | ||
| dummy_functor_3<internal::namespace_short::val_2> f3; | ||
| dummy_functor_4<enum_in_anonNS::val_2> f4; | ||
| dummy_functor_5<no_type_set::val_1> f5; | ||
|
|
||
| cl::sycl::queue q; | ||
|
|
||
| q.submit([&](cl::sycl::handler &cgh) { | ||
| cgh.single_task(f1); | ||
| }); | ||
|
|
||
| q.submit([&](cl::sycl::handler &cgh) { | ||
| cgh.single_task(f2); | ||
| }); | ||
|
|
||
| q.submit([&](cl::sycl::handler &cgh) { | ||
| cgh.single_task(f3); | ||
| }); | ||
|
|
||
| q.submit([&](cl::sycl::handler &cgh) { | ||
| cgh.single_task(f4); | ||
| }); | ||
|
|
||
| q.submit([&](cl::sycl::handler &cgh) { | ||
| cgh.single_task(f5); | ||
| }); | ||
|
|
||
| return 0; | ||
| } | ||
|
|
||
| // CHECK: Forward declarations of templated kernel function types: | ||
| // CHECK: enum class no_namespace_int : int; | ||
| // CHECK: template <no_namespace_int EnumType> class dummy_functor_1; | ||
| // CHECK: enum class no_namespace_short : short; | ||
| // CHECK: template <no_namespace_short EnumType> class dummy_functor_2; | ||
| // CHECK: namespace internal { | ||
| // CHECK-NEXT: enum class namespace_short : short; | ||
| // CHECK-NEXT: } | ||
| // CHECK: template <internal::namespace_short EnumType> class dummy_functor_3; | ||
| // CHECK: namespace { | ||
| // CHECK-NEXT: enum class enum_in_anonNS : short; | ||
| // CHECK-NEXT: } | ||
| // CHECK: template <enum_in_anonNS EnumType> class dummy_functor_4; | ||
| // CHECK: enum class no_type_set : int; | ||
| // CHECK: template <no_type_set EnumType> class dummy_functor_5; | ||
|
|
||
| // CHECK: Specializations of KernelInfo for kernel function types: | ||
| // CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>> | ||
| // CHECK: template <> struct KernelInfo<::dummy_functor_2<(no_namespace_short)1>> | ||
| // CHECK: template <> struct KernelInfo<::dummy_functor_3<(internal::namespace_short)1>> | ||
| // CHECK: template <> struct KernelInfo<::dummy_functor_4<(enum_in_anonNS)1>> | ||
| // CHECK: template <> struct KernelInfo<::dummy_functor_5<(no_type_set)0>> | ||
Uh oh!
There was an error while loading. Please reload this page.