Skip to content
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.

Sign up for GitHub

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

Altering SemaSYCL header gen for wider range of kernel names #46

Closed
wants to merge 1 commit into from
Closed

Altering SemaSYCL header gen for wider range of kernel names #46

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
20 changes: 17 additions & 3 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -736,6 +736,20 @@ static std::string eraseAnonNamespace(std::string S) {
return S;
}

// Removes all "class" and "struct" substrings from given string
static std::string eraseClassAndStruct(std::string S) {
const char S1[] = "class ";
const char S2[] = "struct ";

for (auto Pos = S.find(S1); Pos != StringRef::npos; Pos = S.find(S1, Pos))
S.erase(Pos, sizeof(S1) - 1);

for (auto Pos = S.find(S2); Pos != StringRef::npos; Pos = S.find(S2, Pos))
S.erase(Pos, sizeof(S2) - 1);

return S;
}

// Creates a mangled kernel name for given kernel name type
static std::string constructKernelName(QualType KernelNameType,
ASTContext &AC) {
Expand All @@ -761,8 +775,7 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) {
assert(TemplateArgs && "No template argument info");
// The first template argument always describes the kernel name - whether
// it is lambda or functor.
QualType KernelNameType = TypeName::getFullyQualifiedType(
TemplateArgs->get(0).getAsType(), getASTContext(), true);
QualType KernelNameType = TemplateArgs->get(0).getAsType();
std::string Name = constructKernelName(KernelNameType, getASTContext());
populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE);
FunctionDecl *SYCLKernel =
Expand Down Expand Up @@ -1025,7 +1038,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
for (const KernelDesc &K : KernelDescs) {
const size_t N = K.Params.size();
O << "template <> struct KernelInfo<"
<< eraseAnonNamespace(K.NameType.getAsString()) << "> {\n";
<< eraseClassAndStruct(eraseAnonNamespace(K.NameType.getAsString()))
<< "> {\n";
O << " static constexpr const char* getName() { return \"" << K.Name
<< "\"; }\n";
O << " static constexpr unsigned getNumParams() { return " << N << "; }\n";
Expand Down
31 changes: 28 additions & 3 deletions clang/test/CodeGenSYCL/integration_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,12 +10,17 @@
// CHECK-NEXT: struct X;
// CHECK-NEXT: template <typename T> struct point;
// CHECK-NEXT: template <int a, typename T1, typename T2> class third_kernel;
// CHECK-NEXT: namespace template_arg_ns {
// CHECK-NEXT: template <int DimX> struct namespaced_arg;
// CHECK-NEXT: }
// CHECK-NEXT: template <typename ...Ts> class fourth_kernel;
//
// CHECK: static constexpr
// CHECK-NEXT: const char* const kernel_names[] = {
// CHECK-NEXT: "_ZTSZ4mainE12first_kernel",
// CHECK-NEXT: "_ZTSN16second_namespace13second_kernelIcEE",
// CHECK-NEXT: "_ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE"
// CHECK-NEXT: "_ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE"
// CHECK-NEXT: };
//
// CHECK: static constexpr
Expand All @@ -41,12 +46,19 @@
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 },
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 },
// CHECK-EMPTY:
// CHECK-NEXT: };
//
// CHECK: template <class KernelNameType> struct KernelInfo;
// 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<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> >> {

namespace cl {
namespace sycl {
Expand Down Expand Up @@ -124,6 +136,14 @@ class second_kernel;
template <int a, typename T1, typename T2>
class third_kernel;

namespace template_arg_ns {
template <int DimX>
struct namespaced_arg {};
}

template <typename ...Ts>
class fourth_kernel;

int main() {

cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc1;
Expand Down Expand Up @@ -157,6 +177,11 @@ int main() {
acc2.use();
}
});
kernel_single_task<class fourth_kernel<template_arg_ns::namespaced_arg<1>>>([=]() {
if (i == 13) {
acc2.use();
}
});

return 0;
}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel_functor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,7 +168,7 @@ int main() {
cl::sycl::detail::KernelInfo<Functor1>::getName();
// CHECK: Functor1
cl::sycl::detail::KernelInfo<ns::Functor2>::getName();
// CHECK: ::ns::Functor2
// CHECK: ns::Functor2
cl::sycl::detail::KernelInfo<TmplFunctor<int>>::getName();
// CHECK: TmplFunctor<int>
cl::sycl::detail::KernelInfo<TmplConstFunctor<int>>::getName();
Expand Down
19 changes: 19 additions & 0 deletions sycl/test/regression/kernel_name_class.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,8 @@ namespace nm1 {
namespace nm2 {
class C {};
class KernelName0 : public C {};

template <int X> class KernelName11 {};
} // namespace nm2

class KernelName1;
Expand All @@ -34,8 +36,12 @@ template <> class KernelName3<KernelName1>;
template <> class KernelName4<nm1::nm2::KernelName0> {};
template <> class KernelName4<KernelName1> {};

template <typename... Ts> class KernelName10;

} // namespace nm1

template <typename... Ts> class KernelName12;

static int NumTestCases = 0;

namespace nm3 {
Expand Down Expand Up @@ -233,6 +239,19 @@ struct Wrapper {
});
++NumTestCases;
#endif

deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
cgh.single_task<nm1::KernelName10<nm1::nm2::KernelName11<10>>>([=]() { acc[0] += GOLD; });
});
++NumTestCases;

deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
cgh.single_task<KernelName12<nm1::nm2::KernelName11<10>>>([=]() { acc[0] += GOLD; });
});
++NumTestCases;

}
return arr[0];
}
Expand Down