Skip to content
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
133 changes: 130 additions & 3 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1637,6 +1637,27 @@ static std::string eraseAnonNamespace(std::string S) {
return S;
}

static bool checkIfDeclInCLNameSpace(DeclContext *DC) {

bool inCL = false;

while (DC) {
auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);

if (!NS)
break;

DC = NS->getDeclContext();
if (DC->isTranslationUnit()) {
const IdentifierInfo *II = NS->getIdentifier();
if (II && II->isStr("cl"))
inCL = true;
break;
}
}
return inCL;
}

// Emits a forward declaration
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
SourceLocation KernelLocation) {
Expand Down Expand Up @@ -1690,10 +1711,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();
O << " : " << T.getAsString() << ";\n";
} else
O << ";\n";

// print closing braces for namespaces if needed
for (unsigned I = 0; I < NamespaceCnt; ++I)
Expand Down Expand Up @@ -1796,6 +1824,26 @@ 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>()) {
auto EnumDecl = ET->getDecl();
DeclContext *DC = EnumDecl->getDeclContext();
auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
bool InCLNameSpace = false;

// Header files are included before integration header. Enums
// therefore defined in headers need not be forward declared. Their
// definitions are available in integration header.
if (NS)
InCLNameSpace = checkIfDeclInCLNameSpace(DC);

if (!NS || (NS && !InCLNameSpace))
emitFwdDecl(O, EnumDecl, KernelLocation);
}
break;
}
default:
break; // nop
}
Expand All @@ -1822,6 +1870,84 @@ static std::string getCPPTypeString(QualType Ty) {
return eraseAnonNamespace(Ty.getAsString(P));
}

static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
const PrintingPolicy &P, bool ParameterPack) {

if (!ParameterPack)
ArgOS << "<";

bool FirstArg = true;
const char *Comma = ", ";

for (unsigned I = 0; I < Args.size(); I++) {
const TemplateArgument &Arg = Args[I];
if (Arg.getKind() == TemplateArgument::ArgKind::Pack) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems a lot like a few of the other places we emit arguments based on their kind. Is there a way to dispatch the body of much of this instead?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Maaaaybe. But I think the function body is not similar enough to make this efficient. Both functions are looping over arguments to emit 'something' but what is being emitted is different.

if (Arg.pack_size() && !FirstArg)
ArgOS << Comma;
printArguments(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>();

// Enums defined in SYCL headers are not changed to explicit values since
// the definition is visible in integration header.
if (ET && !checkIfDeclInCLNameSpace(ET->getDecl()->getDeclContext())) {
const llvm::APSInt &Val = Arg.getAsIntegral();
ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
} else {
Arg.print(P, ArgOS);
}
} else {
if (!FirstArg)
ArgOS << Comma;
LangOptions LO;
PrintingPolicy TypePolicy(LO);
TypePolicy.SuppressTypedefs = true;
TypePolicy.SuppressTagKeyword = true;
Arg.print(TypePolicy, ArgOS);
}
FirstArg = false;
}

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);

const TemplateArgumentList &Args = TSD->getTemplateArgs();

// Print template arguments substituting enumerators
printArguments(ArgOS, Args.asArray(), P, false);

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";
Expand Down Expand Up @@ -1938,8 +2064,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
Expand Down
16 changes: 8 additions & 8 deletions clang/test/CodeGenSYCL/int_header1.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4,14 +4,14 @@
// CHECK:template <> struct KernelInfo<class KernelName> {
// CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::nm2::KernelName0>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::KernelName1>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::nm2::KernelName0>> {
// 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>> {
// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::nm2::KernelName0>> {
// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::KernelName1>> {
// CHECK:template <> struct KernelInfo<nm1::KernelName4<nm1::nm2::KernelName0>> {
// 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<TmplClassInAnonNS<ClassInAnonNS>> {

// This test checks if the SYCL device compiler is able to generate correct
// integration header when the kernel name class is expressed in different
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/integration_header.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -56,9 +56,9 @@
// CHECK-NEXT: };
//
// 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> >> {
// 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"

Expand Down
30 changes: 15 additions & 15 deletions clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -100,37 +100,37 @@ struct kernel_name4;
int main() {
dummy_functor f;
// non-type template arguments
// CHECK: template <> struct KernelInfo<::kernel_name1<1, 1>> {
// CHECK: template <> struct KernelInfo<kernel_name1<1, 1>> {
single_task<kernel_name1<1, 1>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name1v1<1, 1>> {
// CHECK: template <> struct KernelInfo<kernel_name1v1<1, 1>> {
single_task<kernel_name1v1<1, 1>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name1v2<1, 1>> {
// CHECK: template <> struct KernelInfo<kernel_name1v2<1, 1>> {
single_task<kernel_name1v2<1, 1>>(f);
// partial template specialization
// CHECK: template <> struct KernelInfo<::kernel_name2<int, int>> {
// CHECK: template <> struct KernelInfo<kernel_name2<int, int>> {
single_task<kernel_name2<int_t, int>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2<const int, char>> {
// CHECK: template <> struct KernelInfo<kernel_name2<const int, char>> {
single_task<kernel_name2<cint_t, char>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2<long, float>> {
// CHECK: template <> struct KernelInfo<kernel_name2<long, float>> {
single_task<kernel_name2<space::long_t, float>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2<const long, ::A>> {
// CHECK: template <> struct KernelInfo<kernel_name2<const long, A>> {
single_task<kernel_name2<const space::long_t, space::a_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2<const volatile long, const ::space::B>> {
// CHECK: template <> struct KernelInfo<kernel_name2<const volatile long, const space::B>> {
single_task<kernel_name2<volatile space::clong_t, const space::b_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2< ::A, long>> {
// CHECK: template <> struct KernelInfo<kernel_name2<A, long>> {
single_task<kernel_name2<space::a_t, space::long_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2< ::space::B, int>> {
// CHECK: template <> struct KernelInfo<kernel_name2<space::B, int>> {
single_task<kernel_name2<space::b_t, int_t>>(f);
// full template specialization
// CHECK: template <> struct KernelInfo<::kernel_name2<int, const unsigned int>> {
// CHECK: template <> struct KernelInfo<kernel_name2<int, const unsigned int>> {
single_task<kernel_name2<int_t, const uint_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2<const long, volatile const unsigned long>> {
// CHECK: template <> struct KernelInfo<kernel_name2<const long, const volatile unsigned long>> {
single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name2< ::A, volatile ::space::B>> {
// CHECK: template <> struct KernelInfo<kernel_name2<A, volatile space::B>> {
single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name3<1>> {
// CHECK: template <> struct KernelInfo<kernel_name3<1>> {
single_task<kernel_name3<1>>(f);
// CHECK: template <> struct KernelInfo<::kernel_name4<1>> {
// CHECK: template <> struct KernelInfo<kernel_name4<1>> {
single_task<kernel_name4<1>>(f);

return 0;
Expand Down
100 changes: 100 additions & 0 deletions clang/test/CodeGenSYCL/kernelname-enum.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,100 @@
// 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 {
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
};
}

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()() {}
};

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;

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);
});

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: 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>>