From de57011f53c9b1abe3ac2efc1503c596f74554c1 Mon Sep 17 00:00:00 2001 From: Andrew Gozillon Date: Thu, 21 Mar 2019 14:53:38 -0700 Subject: [PATCH] Altering SemaSYCL header gen for wider range of kernel names The goal of this patch was to allow kernel names of the type: Example 1: // kernel name components namespace nm1 { namespace nm2 { template class K0 {}; } template class K1; } // The type to be the kernel name nm1::K1> Example 2: // kernel name components nm { template class K0 {}; } template class K1; // The type to be the kernel name K1> To do this I altered SemaSYCL.cpp to no longer use TypeName::getFullyQualifiedType when assigning QualType KernelNameType and instead just use the qualified type that is retrieved from the TemplateArgs itself. This seems to work. Passing all previous tests and failing the previous tests that it was failing before, whilst passing the two new additional cases I looked at. I may be missing some context as to why getFullyQualifiedType was used previously (it seems less qualified than the actual return from getAsType, missing namespace qualifiers on template arguments) and why there was a reliance on the :: global scope resolution operator i.e. ::type. So perhaps I am missing some information like an important use case or implementation detail that may make this an invalid change. I also added another string alteration function that runs over the qualified name and removes class/struct where relevant in the specialization. This isn't needed for the newly tested way of naming a kernel, but the generated header seems a little more correct in the sense that the header specializations are no longer redefining the class/struct inline on the specialization and are instead reusing the definitions at the top of the integration header. I made modifcations to the integration_header.cpp, kernel_functor.cpp and kernel_name_class.cpp. I changed integration_header.cpp and kernel_functor.cpp checks to test for the new/changed output of the integration header generation. I added new test cases to integration_header.cpp and kernel_name_class.cpp that reflect the test cases this SemaSYCL.cpp modification aims to fix. Signed-off-by: Andrew Gozillon --- clang/lib/Sema/SemaSYCL.cpp | 20 ++++++++++-- clang/test/CodeGenSYCL/integration_header.cpp | 31 +++++++++++++++++-- clang/test/CodeGenSYCL/kernel_functor.cpp | 2 +- sycl/test/regression/kernel_name_class.cpp | 19 ++++++++++++ 4 files changed, 65 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c7301d5970b6e..789da27fc0091 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -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) { @@ -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 = @@ -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"; diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 9e457819562d9..cf46052437a46 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -10,12 +10,17 @@ // CHECK-NEXT: struct X; // CHECK-NEXT: template struct point; // CHECK-NEXT: template class third_kernel; +// CHECK-NEXT: namespace template_arg_ns { +// CHECK-NEXT: template struct namespaced_arg; +// CHECK-NEXT: } +// CHECK-NEXT: template 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 @@ -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 struct KernelInfo; -// CHECK: template <> struct KernelInfo { -// CHECK: template <> struct KernelInfo<::second_namespace::second_kernel> { -// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point >> { +// CHECK: template <> struct KernelInfo { +// CHECK: template <> struct KernelInfo> { +// CHECK: template <> struct KernelInfo >> { +// CHECK: template <> struct KernelInfo >> { namespace cl { namespace sycl { @@ -124,6 +136,14 @@ class second_kernel; template class third_kernel; +namespace template_arg_ns { + template + struct namespaced_arg {}; +} + +template +class fourth_kernel; + int main() { cl::sycl::accessor acc1; @@ -157,6 +177,11 @@ int main() { acc2.use(); } }); + kernel_single_task>>([=]() { + if (i == 13) { + acc2.use(); + } + }); return 0; } diff --git a/clang/test/CodeGenSYCL/kernel_functor.cpp b/clang/test/CodeGenSYCL/kernel_functor.cpp index 56b9b34c2dbeb..c37f6a343b1de 100644 --- a/clang/test/CodeGenSYCL/kernel_functor.cpp +++ b/clang/test/CodeGenSYCL/kernel_functor.cpp @@ -168,7 +168,7 @@ int main() { cl::sycl::detail::KernelInfo::getName(); // CHECK: Functor1 cl::sycl::detail::KernelInfo::getName(); - // CHECK: ::ns::Functor2 + // CHECK: ns::Functor2 cl::sycl::detail::KernelInfo>::getName(); // CHECK: TmplFunctor cl::sycl::detail::KernelInfo>::getName(); diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp index 4e4d2e835ef55..9de14e1ba99e6 100644 --- a/sycl/test/regression/kernel_name_class.cpp +++ b/sycl/test/regression/kernel_name_class.cpp @@ -21,6 +21,8 @@ namespace nm1 { namespace nm2 { class C {}; class KernelName0 : public C {}; + +template class KernelName11 {}; } // namespace nm2 class KernelName1; @@ -34,8 +36,12 @@ template <> class KernelName3; template <> class KernelName4 {}; template <> class KernelName4 {}; +template class KernelName10; + } // namespace nm1 +template class KernelName12; + static int NumTestCases = 0; namespace nm3 { @@ -233,6 +239,19 @@ struct Wrapper { }); ++NumTestCases; #endif + + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task>>([=]() { acc[0] += GOLD; }); + }); + ++NumTestCases; + + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task>>([=]() { acc[0] += GOLD; }); + }); + ++NumTestCases; + } return arr[0]; }