diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7cd31c670852f..c351fd343fe9b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4491,10 +4491,22 @@ SYCLIntegrationHeader::SYCLIntegrationHeader(bool _UnnamedLambdaSupport, : UnnamedLambdaSupport(_UnnamedLambdaSupport), S(_S) {} void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) { + // Skip the dependent version of these variables, we only care about them + // after instantiation. + if (VD->getDeclContext()->isDependentContext()) + return; // Step 1: ensure that this is of the correct type-spec-constant template // specialization). - if (!Util::isSyclSpecIdType(VD->getType())) - return; + if (!Util::isSyclSpecIdType(VD->getType())) { + // Handle the case where this could be a deduced type, such as a deduction + // guide. We have to do this here since this function, unlike most of the + // rest of this file, is called during Sema instead of after it. We will + // also have to filter out after deduction later. + QualType Ty = VD->getType().getCanonicalType(); + + if (!Ty->isUndeducedType()) + return; + } // Step 2: ensure that this is a static member, or a namespace-scope. // Note that isLocalVarDeclorParm excludes thread-local and static-local // intentionally, as there is no way to 'spell' one of those in the @@ -4536,26 +4548,162 @@ void SYCLIntegrationFooter::emitSpecIDName(raw_ostream &O, const VarDecl *VD) { O << ""; } -bool SYCLIntegrationFooter::emit(raw_ostream &O) { +template +static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS, + const DeclContext *DC) { + if (DC->isTranslationUnit()) + return; + + const auto *CurDecl = cast(DC); + // Ensure we are in the canonical version, so that we know we have the 'full' + // name of the thing. + CurDecl = CurDecl->getCanonicalDecl(); + + // We are intentionally skipping linkage decls and record decls. Namespaces + // can appear in a linkage decl, but not a record decl, so we don't have to + // worry about the names getting messed up from that. We handle record decls + // later when printing the name of the thing. + const auto *NS = dyn_cast(CurDecl); + if (NS) + Before(OS, NS); + + if (const DeclContext *NewDC = CurDecl->getDeclContext()) + PrintNSHelper(Before, After, OS, NewDC); + + if (NS) + After(OS, NS); +} + +static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC) { + PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {}, + [](raw_ostream &OS, const NamespaceDecl *NS) { + if (NS->isInline()) + OS << "inline "; + OS << "namespace "; + if (!NS->isAnonymousNamespace()) + OS << NS->getName() << " "; + OS << "{\n"; + }, + OS, DC); +} + +static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { + PrintNSHelper( + [](raw_ostream &OS, const NamespaceDecl *NS) { + OS << "} // "; + if (NS->isInline()) + OS << "inline "; + + OS << "namespace "; + if (!NS->isAnonymousNamespace()) + OS << NS->getName(); + + OS << '\n'; + }, + [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC); +} + +static std::string EmitSpecIdShim(raw_ostream &OS, unsigned &ShimCounter, + const std::string &LastShim, + const NamespaceDecl *AnonNS) { + std::string NewShimName = + "__sycl_detail::__spec_id_shim_" + std::to_string(ShimCounter) + "()"; + // Print opening-namespace + PrintNamespaces(OS, Decl::castToDeclContext(AnonNS)); + OS << "namespace __sycl_detail {\n"; + OS << "static constexpr decltype(" << LastShim << ") &__spec_id_shim_" + << ShimCounter << "() {\n"; + OS << " return " << LastShim << ";\n"; + OS << "}\n"; + OS << "} // namespace __sycl_detail \n"; + PrintNSClosingBraces(OS, Decl::castToDeclContext(AnonNS)); + + ++ShimCounter; + return std::move(NewShimName); +} + +// Emit the list of shims required for a DeclContext, calls itself recursively. +static void EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter, + const DeclContext *DC, + std::string &NameForLastShim) { + if (DC->isTranslationUnit()) { + NameForLastShim = "::" + NameForLastShim; + return; + } + + const auto *CurDecl = cast(DC)->getCanonicalDecl(); + + // We skip linkage decls, since they don't modify the Qualified name. + if (const auto *RD = dyn_cast(CurDecl)) { + NameForLastShim = RD->getNameAsString() + "::" + NameForLastShim; + } else if (const auto *ND = dyn_cast(CurDecl)) { + if (ND->isAnonymousNamespace()) { + // Print current shim, reset 'name for last shim'. + NameForLastShim = EmitSpecIdShim(OS, ShimCounter, NameForLastShim, ND); + } else { + NameForLastShim = ND->getNameAsString() + "::" + NameForLastShim; + } + } else { + // FIXME: I don't believe there are other declarations that these variables + // could possibly find themselves in. LinkageDecls don't change the + // qualified name, so there is nothing to do here. At one point we should + // probably convince ourselves that this is entire list and remove this + // comment. + assert((isa(CurDecl)) && + "Unhandled decl type"); + } + + EmitSpecIdShims(OS, ShimCounter, CurDecl->getDeclContext(), NameForLastShim); +} + +// Emit the list of shims required for a variable declaration. +// Returns a string containing the FQN of the 'top most' shim, including its +// function call parameters. +static std::string EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter, + const VarDecl *VD) { + assert(VD->isInAnonymousNamespace() && + "Function assumes this is in an anonymous namespace"); + std::string RelativeName = VD->getNameAsString(); + EmitSpecIdShims(OS, ShimCounter, VD->getDeclContext(), RelativeName); + return std::move(RelativeName); +} + +bool SYCLIntegrationFooter::emit(raw_ostream &OS) { PrintingPolicy Policy{S.getLangOpts()}; Policy.adjustForCPlusPlusFwdDecl(); Policy.SuppressTypedefs = true; Policy.SuppressUnwrittenScope = true; - for (const VarDecl *D : SpecConstants) { - O << "template<>\n"; - O << "inline const char *get_spec_constant_symbolic_ID<"; - // Emit the FQN for this, but we probably need to do some funny-business for - // anonymous namespaces. - D->printQualifiedName(O, Policy); - O << ">() {\n"; - O << " return \""; - emitSpecIDName(O, D); - O << "\";\n"; - O << "}\n"; + // Used to uniquely name the 'shim's as we generate the names in each + // anonymous namespace. + unsigned ShimCounter = 0; + for (const VarDecl *VD : SpecConstants) { + VD = VD->getCanonicalDecl(); + if (VD->isInAnonymousNamespace()) { + std::string TopShim = EmitSpecIdShims(OS, ShimCounter, VD); + OS << "namespace sycl {\n"; + OS << "namespace detail {\n"; + OS << "template<>\n"; + OS << "inline const char *get_spec_constant_symbolic_ID<" << TopShim + << ">() {\n"; + OS << " return " << TopShim << ";\n"; + } else { + OS << "namespace sycl {\n"; + OS << "namespace detail {\n"; + OS << "template<>\n"; + OS << "inline const char *get_spec_constant_symbolic_ID<::"; + VD->printQualifiedName(OS, Policy); + OS << ">() {\n"; + OS << " return \""; + emitSpecIDName(OS, VD); + OS << "\";\n"; + } + OS << "}\n"; + OS << "} // namespace detail\n"; + OS << "} // namespace sycl\n"; } - O << "#include \n"; + OS << "#include \n"; return true; } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 6f0b1281b86da..d27f3c9e1a67d 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -321,6 +321,10 @@ template class specialization_id { T MDefaultValue; }; +#if __cplusplus >= 201703L +template specialization_id(T) -> specialization_id; +#endif // C++17. + #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTask diff --git a/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp b/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp new file mode 100644 index 0000000000000..ccb216dc85b53 --- /dev/null +++ b/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp @@ -0,0 +1,363 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll +// RUN: FileCheck -input-file=%t.h %s +// A test that validates the more complex cases of the specialization-constant +// integration footer details, basically any situation we can come up with that +// has an anonymous namespace. + +#include "Inputs/sycl.hpp" +int main() { + cl::sycl::kernel_single_task([]() {}); +} + +using namespace cl; +// Example ways in which the application can declare a "specialization_id" +// variable. +struct S1 { + static constexpr sycl::specialization_id a{1}; + // CHECK: namespace sycl { + // CHECK-NEXT: namespace detail { + // CHECK-NEXT: template<> + // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() { + // CHECK-NEXT: return ""; + // CHECK-NEXT: } + // CHECK-NEXT: } // namespace detail + // CHECK-NEXT: } // namespace sycl +}; +constexpr sycl::specialization_id b{2}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +inline constexpr sycl::specialization_id c{3}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +static constexpr sycl::specialization_id d{4}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl + +namespace { +struct S2 { + static constexpr sycl::specialization_id a{18}; + // CHECK-NEXT: namespace { + // CHECK-NEXT: namespace __sycl_detail { + // CHECK-NEXT: static constexpr decltype(S2::a) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { + // CHECK-NEXT: return S2::a; + // CHECK-NEXT: } + // CHECK-NEXT: } // namespace __sycl_detail + // CHECK-NEXT: } // namespace + // CHECK-NEXT: namespace sycl { + // CHECK-NEXT: namespace detail { + // CHECK-NEXT: template<> + // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { + // CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); + // CHECK-NEXT: } + // CHECK-NEXT: // namespace detail + // CHECK-NEXT: // namespace sycl +}; +} // namespace + +template +struct S3 { + static constexpr sycl::specialization_id a{Val}; +}; +template class S3<1>; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<1>::a>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +template class S3<2>; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S3<2>::a>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl + +namespace inner { +constexpr sycl::specialization_id same_name{5}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +} +constexpr sycl::specialization_id same_name{6}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +namespace { +constexpr sycl::specialization_id same_name{7}; +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +namespace { +namespace inner { +constexpr sycl::specialization_id same_name{8}; +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(inner::same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return inner::same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} // namespace +namespace inner { +namespace { +constexpr sycl::specialization_id same_name{9}; +// CHECK-NEXT: namespace inner { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace inner +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} // namespace inner + +namespace outer { +constexpr sycl::specialization_id same_name{10}; +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +namespace { +constexpr sycl::specialization_id same_name{11}; +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl + +namespace inner { +constexpr sycl::specialization_id same_name{12}; +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(inner::same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return inner::same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl + +namespace { +// This has multiple anonymous namespaces in its declaration context, we need to +// make sure we emit a shim for each. +constexpr sycl::specialization_id same_name{13}; +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace inner { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace inner +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer + +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()) &__spec_id_shim_[[SHIM_ID_2:[0-9]+]]() { +// CHECK-NEXT: return inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() { +// CHECK-NEXT: return ::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} // namespace inner +} // namespace +} // namespace outer + +namespace { +namespace outer { +constexpr sycl::specialization_id same_name{14}; +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(outer::same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return outer::same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +namespace { +constexpr sycl::specialization_id same_name{15}; +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()) &__spec_id_shim_[[SHIM_ID2:[0-9]+]]() { +// CHECK-NEXT: return outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace + +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID2]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +namespace inner { +constexpr sycl::specialization_id same_name{16}; +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace outer { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(inner::same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return inner::same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()) &__spec_id_shim_[[SHIM_ID2:[0-9]+]]() { +// CHECK-NEXT: return outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID2]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} // namespace +} // namespace outer +} // namespace + +namespace outer { +namespace inner { +constexpr sycl::specialization_id same_name{17}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl +} +} // namespace outer + +// CHECK: #include diff --git a/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp b/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp new file mode 100644 index 0000000000000..0aa28fcbec3ce --- /dev/null +++ b/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp @@ -0,0 +1,137 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll +// RUN: FileCheck -input-file=%t.h %s +// A test that validates the more complex cases of the specialization-constant +// integration footer details, basically any situation we can come up with that +// has an anonymous namespace. + +#include "Inputs/sycl.hpp" +int main() { + cl::sycl::kernel_single_task([]() {}); +} +using namespace cl; + +struct S1 { + static constexpr sycl::specialization_id a{1}; +}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::S1::a>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl + +constexpr sycl::specialization_id b{202}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::b>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +inline constexpr sycl::specialization_id c{3}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::c>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +static constexpr sycl::specialization_id d{205}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::d>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl + +namespace inner { +constexpr sycl::specialization_id same_name{5}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::inner::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +constexpr sycl::specialization_id same_name{6}; +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::same_name>() { +// CHECK-NEXT: return ""; +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +namespace { +constexpr sycl::specialization_id same_name{207}; +// CHECK: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +namespace { +namespace inner { +constexpr sycl::specialization_id same_name{208}; +// CHECK: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(inner::same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return inner::same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} // namespace + +namespace outer::inline inner { +namespace { +constexpr sycl::specialization_id same_name{209}; +// CHECK: namespace outer { +// CHECK-NEXT: namespace inner { +// CHECK-NEXT: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(same_name) &__spec_id_shim_[[SHIM_ID:[0-9]+]]() { +// CHECK-NEXT: return same_name; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // inline namespace inner +// CHECK-NEXT: } // namespace outer +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() { +// CHECK-NEXT: return ::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]](); +// CHECK-NEXT: } +// CHECK-NEXT: // namespace detail +// CHECK-NEXT: // namespace sycl +} +} + +// CHECK: #include diff --git a/clang/test/CodeGenSYCL/integration_footer.cpp b/clang/test/CodeGenSYCL/integration_footer.cpp index 41b37f7b5fd68..b1e4846537cb1 100644 --- a/clang/test/CodeGenSYCL/integration_footer.cpp +++ b/clang/test/CodeGenSYCL/integration_footer.cpp @@ -10,17 +10,25 @@ int main() { using namespace cl::sycl; cl::sycl::specialization_id GlobalSpecID; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::GlobalSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl struct Wrapper { static specialization_id WrapperSpecID; - // CHECK: template<> - // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { + // CHECK: namespace sycl { + // CHECK-NEXT: namespace detail { + // CHECK-NEXT: template<> + // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Wrapper::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } + // CHECK-NEXT: } // namespace detail + // CHECK-NEXT: } // namespace sycl }; template @@ -28,40 +36,64 @@ struct WrapperTemplate { static specialization_id WrapperSpecID; }; template class WrapperTemplate; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID::WrapperSpecID>() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::WrapperTemplate::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl template class WrapperTemplate; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID::WrapperSpecID>() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::WrapperTemplate::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl namespace Foo { specialization_id NSSpecID; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::NSSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl inline namespace Bar { specialization_id InlineNSSpecID; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::InlineNSSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl specialization_id NSSpecID; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::Bar::NSSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl struct Wrapper { static specialization_id WrapperSpecID; - // CHECK: template<> - // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { + // CHECK: namespace sycl { + // CHECK-NEXT: namespace detail { + // CHECK-NEXT: template<> + // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::Wrapper::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } + // CHECK-NEXT: } // namespace detail + // CHECK-NEXT: } // namespace sycl }; template @@ -69,22 +101,45 @@ struct WrapperTemplate { static specialization_id WrapperSpecID; }; template class WrapperTemplate; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID::WrapperSpecID>() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::WrapperTemplate::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl template class WrapperTemplate; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID::WrapperSpecID>() { +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::WrapperTemplate::WrapperSpecID>() { // CHECK-NEXT: return ""; // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl } // namespace Bar namespace { specialization_id AnonNSSpecID; -// CHECK: template<> -// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID() { -// CHECK-NEXT: return ""; + +// CHECK: namespace Foo { +// CHECK: namespace { +// CHECK-NEXT: namespace __sycl_detail { +// CHECK-NEXT: static constexpr decltype(AnonNSSpecID) &__spec_id_shim_[[SHIM0:[0-9]+]]() { +// CHECK-NEXT: return AnonNSSpecID; +// CHECK-NEXT: } +// CHECK-NEXT: } // namespace __sycl_detail +// CHECK-NEXT: } // namespace +// CHECK-NEXT: } // namespace Foo + +// CHECK: namespace sycl { +// CHECK-NEXT: namespace detail { +// CHECK-NEXT: template<> +// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID<::Foo::__sycl_detail::__spec_id_shim_[[SHIM0]]()>() { +// CHECK-NEXT: return ::Foo::__sycl_detail::__spec_id_shim_[[SHIM0]](); // CHECK-NEXT: } +// CHECK-NEXT: } // namespace detail +// CHECK-NEXT: } // namespace sycl } // namespace } // namespace Foo