From 811239578d3c8ceeb518eba8f688c2223692fdda Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 11:44:18 +0300
Subject: [PATCH 01/11] [SYCL] Emit integration footer and header for
 device_global variables

The implementation is based on the design doc
https://github.com/intel/llvm/blob/sycl/sycl/doc/DeviceGlobal.md .
---
 clang/include/clang/Basic/Attr.td             |   8 +
 clang/include/clang/Sema/Sema.h               |   4 +-
 clang/lib/Sema/Sema.cpp                       |   7 +-
 clang/lib/Sema/SemaDeclAttr.cpp               |   3 +
 clang/lib/Sema/SemaSYCL.cpp                   | 139 +++++++++++++-----
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        |  10 ++
 .../anonymous_integration_footer.cpp          |  52 +++----
 .../anonymous_integration_footer2.cpp         |  12 +-
 .../device_global_int_footer_header.cpp       | 109 ++++++++++++++
 clang/test/CodeGenSYCL/integration_footer.cpp |   8 +-
 10 files changed, 279 insertions(+), 73 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/device_global_int_footer_header.cpp

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 3b23bdfc3c130..e963e64661eb5 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1261,6 +1261,14 @@ def SYCLUsesAspects : InheritableAttr {
   let Documentation = [Undocumented];
 }
 
+def SYCLDeviceGlobal : InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "device_global">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  // Only used internally by the SYCL implementation
+  let Documentation = [Undocumented];
+}
+
 // Marks functions which must not be vectorized via horizontal SIMT widening,
 // e.g. because the function is already vectorized. Used to mark SYCL
 // explicit SIMD kernels and functions.
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 133d9d0f5b715..d3091268b2634 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -440,12 +440,14 @@ class SYCLIntegrationFooter {
   SYCLIntegrationFooter(Sema &S) : S(S) {}
   bool emit(StringRef MainSrc);
   void addVarDecl(const VarDecl *VD);
+  bool metSYCLDeviceGlobals() { return DeviceGlobalsEmitted; }
 
 private:
   bool emit(raw_ostream &O);
   Sema &S;
-  llvm::SmallVector<const VarDecl *> SpecConstants;
+  llvm::SmallVector<const VarDecl *> GlobalVars;
   void emitSpecIDName(raw_ostream &O, const VarDecl *VD);
+  bool DeviceGlobalsEmitted = false;
 };
 
 /// Tracks expected type during expression parsing, for use in code completion.
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index c82d79339643d..115290145d8a4 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1101,11 +1101,14 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
     // Set the names of the kernels, now that the names have settled down. This
     // needs to happen before we generate the integration headers.
     SetSYCLKernelNames();
+    // Make sure that the footer emitted before header, since only after the
+    // footer is emitted it is known that translation unit contains device
+    // global variables.
+    if (SyclIntFooter != nullptr)
+      SyclIntFooter->emit(getLangOpts().SYCLIntFooter);
     // Emit SYCL integration header for current translation unit if needed
     if (SyclIntHeader != nullptr)
       SyclIntHeader->emit(getLangOpts().SYCLIntHeader);
-    if (SyclIntFooter != nullptr)
-      SyclIntFooter->emit(getLangOpts().SYCLIntFooter);
     MarkDevices();
   }
 
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 186cdf508fcad..3936b230e1e98 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -10414,6 +10414,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLSpecialClass:
     handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLDeviceGlobal:
+    handleSimpleAttribute<SYCLDeviceGlobalAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLDevice:
     handleSYCLDeviceAttr(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index cfd95a9085367..92d5901c49b71 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -123,6 +123,10 @@ class Util {
   /// specialization id class.
   static bool isSyclSpecIdType(QualType Ty);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// device_global class.
+  static bool isSyclDeviceGlobalType(QualType Ty);
+
   /// Checks whether given clang type is a full specialization of the SYCL
   /// kernel_handler class.
   static bool isSyclKernelHandlerType(QualType Ty);
@@ -4676,7 +4680,23 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "namespace sycl {\n";
   O << "namespace detail {\n";
 
-  O << "\n";
+  // Generate declaration of variable of type __sycl_device_global_registration
+  // whose sole purpose is to run its constructor before the application's
+  // main() function.
+
+  if (S.getSyclIntegrationFooter().metSYCLDeviceGlobals()) {
+    O << "namespace {\n";
+
+    O << "class __sycl_device_global_registration {\n";
+    O << "public:\n";
+    O << "  __sycl_device_global_registration() noexcept;\n";
+    O << "};\n";
+    O << "__sycl_device_global_registration __sycl_device_global_registerer;\n";
+
+    O << "} // namespace\n";
+
+    O << "\n";
+  }
 
   O << "// names of all kernels defined in the corresponding source\n";
   O << "static constexpr\n";
@@ -4858,9 +4878,9 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
   // template instantiations as a VarDecl.
   if (isa<VarTemplatePartialSpecializationDecl>(VD))
     return;
-  // Step 1: ensure that this is of the correct type-spec-constant template
-  // specialization).
-  if (!Util::isSyclSpecIdType(VD->getType())) {
+  // Step 1: ensure that this is of the correct type template specialization.
+  if (!Util::isSyclSpecIdType(VD->getType()) &&
+      !Util::isSyclDeviceGlobalType(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
@@ -4876,8 +4896,8 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
   // let an error happen during host compilation.
   if (!VD->hasGlobalStorage() || VD->isLocalVarDeclOrParm())
     return;
-  // Step 3: Add to SpecConstants collection.
-  SpecConstants.push_back(VD);
+  // Step 3: Add to collection.
+  GlobalVars.push_back(VD);
 }
 
 // Post-compile integration header support.
@@ -4955,15 +4975,15 @@ 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) + "()";
+      "__sycl_detail::__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_"
+  OS << "static constexpr decltype(" << LastShim << ") &__shim_"
      << ShimCounter << "() {\n";
   OS << "  return " << LastShim << ";\n";
   OS << "}\n";
-  OS << "} // namespace __sycl_detail \n";
+  OS << "} // namespace __sycl_detail\n";
   PrintNSClosingBraces(OS, Decl::castToDeclContext(AnonNS));
 
   ++ShimCounter;
@@ -5026,58 +5046,97 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   Policy.SuppressTypedefs = true;
   Policy.SuppressUnwrittenScope = true;
 
-  llvm::SmallSet<const VarDecl *, 8> VisitedSpecConstants;
+  llvm::SmallSet<const VarDecl *, 8> Visited;
   bool EmittedFirstSpecConstant = false;
 
   // Used to uniquely name the 'shim's as we generate the names in each
   // anonymous namespace.
   unsigned ShimCounter = 0;
-  for (const VarDecl *VD : SpecConstants) {
+
+  std::string DeviceGlobalsBuf;
+  llvm::raw_string_ostream DeviceGlobOS(DeviceGlobalsBuf);
+  for (const VarDecl *VD : GlobalVars) {
     VD = VD->getCanonicalDecl();
 
-    // Skip if this isn't a SpecIdType.  This can happen if it was a deduced
-    // type.
-    if (!Util::isSyclSpecIdType(VD->getType()))
+    // Skip if this isn't a SpecIdType or DeviceGlobal.  This can happen if it
+    // was a deduced type.
+    if (!Util::isSyclSpecIdType(VD->getType()) &&
+        !Util::isSyclDeviceGlobalType(VD->getType()))
       continue;
 
     // Skip if we've already visited this.
-    if (llvm::find(VisitedSpecConstants, VD) != VisitedSpecConstants.end())
+    if (llvm::find(Visited, VD) != Visited.end())
       continue;
 
     // We only want to emit the #includes if we have a spec-constant that needs
     // them, so emit this one on the first time through the loop.
-    if (!EmittedFirstSpecConstant)
+    if (!EmittedFirstSpecConstant && !DeviceGlobalsEmitted)
       OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
-    EmittedFirstSpecConstant = true;
 
-    VisitedSpecConstants.insert(VD);
+
+    Visited.insert(VD);
     std::string TopShim = EmitSpecIdShims(OS, ShimCounter, Policy, VD);
-    OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
-    OS << "namespace sycl {\n";
-    OS << "namespace detail {\n";
-    OS << "template<>\n";
-    OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
+    if (Util::isSyclDeviceGlobalType(VD->getType())) {
+      if (!DeviceGlobalsEmitted)
+        OS << "#include <CL/sycl/detail/device_global_map.hpp>\n";
+
+      DeviceGlobalsEmitted = true;
+      DeviceGlobOS << "device_global_map::add(";
+    } else {
+      EmittedFirstSpecConstant = true;
+      OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
+      OS << "namespace sycl {\n";
+      OS << "namespace detail {\n";
+      OS << "template<>\n";
+      OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
+    }
 
+    std::string VarRefName;
+    llvm::raw_string_ostream VarRefNameOS(VarRefName);
     if (VD->isInAnonymousNamespace()) {
-      OS << TopShim;
+      VarRefNameOS << TopShim;
+    } else {
+      VarRefNameOS << "::";
+      VD->getNameForDiagnostic(VarRefNameOS, Policy, true);
+    }
+    VarRefNameOS.flush();
+    if (Util::isSyclDeviceGlobalType(VD->getType())) {
+      DeviceGlobOS << "&";
+      DeviceGlobOS << VarRefName;
     } else {
-      OS << "::";
-      VD->getNameForDiagnostic(OS, Policy, true);
+      OS << VarRefName;
     }
 
-    OS << ">() {\n";
-    OS << "  return \"";
-    OS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD);
-    OS << "\";\n";
-    OS << "}\n";
-    OS << "} // namespace detail\n";
-    OS << "} // namespace sycl\n";
-    OS << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
+    if (Util::isSyclDeviceGlobalType(VD->getType())) {
+      DeviceGlobOS << ", \"";
+      DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD);
+      DeviceGlobOS << "\");\n";
+    } else {
+      OS << ">() {\n";
+      OS << "  return \"";
+      OS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD);
+      OS << "\";\n";
+      OS << "}\n";
+      OS << "} // namespace detail\n";
+      OS << "} // namespace sycl\n";
+      OS << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
+    }
   }
 
   if (EmittedFirstSpecConstant)
     OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
 
+  if (DeviceGlobalsEmitted) {
+    DeviceGlobOS.flush();
+    OS << "namespace sycl::detail {\n";
+    OS << "namespace {\n";
+    OS << "__sycl_device_global_registration::__sycl_device_global_"
+          "registration() noexcept {\n";
+    OS << DeviceGlobalsBuf;
+    OS << "}\n";
+    OS << "} // namespace (unnamed)\n";
+    OS << "} // namespace sycl::detail\n";
+  }
   return true;
 }
 
@@ -5122,6 +5181,18 @@ bool Util::isSyclSpecIdType(QualType Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclDeviceGlobalType(QualType Ty) {
+  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+  if (!RecTy)
+    return false;
+  if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecTy)) {
+    ClassTemplateDecl *Template = CTSD->getSpecializedTemplate();
+    if (CXXRecordDecl *RD = Template->getTemplatedDecl())
+        return RD->hasAttr<SYCLDeviceGlobalAttr>();
+  }
+  return RecTy->hasAttr<SYCLDeviceGlobalAttr>();
+}
+
 bool Util::isSyclKernelHandlerType(QualType Ty) {
   std::array<DeclContextDesc, 3> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index a883c8a41cc16..465d97b20bcce 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -129,6 +129,16 @@ struct no_alias {
   template <bool> class instance {};
 };
 } // namespace property
+
+template <typename T>
+class [[__sycl_detail__::device_global]] device_global {
+public:
+  const T& get() const noexcept { return *Data; }
+  device_global() = default;
+private:
+  T *Data;
+};
+
 } // namespace oneapi
 } // namespace ext
 
diff --git a/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp b/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp
index 2b5e2ffbf3285..dacb8ca20b608 100644
--- a/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp
+++ b/clang/test/CodeGenSYCL/anonymous_integration_footer.cpp
@@ -67,7 +67,7 @@ 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: static constexpr decltype(S2::a) &__shim_[[SHIM_ID:[0-9]+]]() {
   // CHECK-NEXT: return S2::a;
   // CHECK-NEXT: }
   // CHECK-NEXT: } // namespace __sycl_detail
@@ -76,7 +76,7 @@ struct S2 {
   // CHECK-NEXT: namespace sycl {
   // CHECK-NEXT: namespace detail {
   // CHECK-NEXT: template<>
-  // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+  // CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
   // CHECK-NEXT: return "____ZN12_GLOBAL__N_12S21aE";
   // CHECK-NEXT: }
   // CHECK-NEXT: } // namespace detail
@@ -140,7 +140,7 @@ 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -149,7 +149,7 @@ constexpr sycl::specialization_id same_name{7};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -161,7 +161,7 @@ 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: static constexpr decltype(inner::same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return inner::same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -170,7 +170,7 @@ constexpr sycl::specialization_id same_name{8};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -184,7 +184,7 @@ 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -194,7 +194,7 @@ constexpr sycl::specialization_id same_name{9};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::inner::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN5inner12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -220,7 +220,7 @@ 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -230,7 +230,7 @@ constexpr sycl::specialization_id same_name{11};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -242,7 +242,7 @@ 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: static constexpr decltype(inner::same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return inner::same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -252,7 +252,7 @@ constexpr sycl::specialization_id same_name{12};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner9same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -268,7 +268,7 @@ constexpr sycl::specialization_id same_name{13};
 // 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -280,8 +280,8 @@ constexpr sycl::specialization_id same_name{13};
 // 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: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM_ID]]()) &__shim_[[SHIM_ID_2:[0-9]+]]() {
+// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM_ID]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
 // CHECK-NEXT: } // namespace
@@ -290,7 +290,7 @@ constexpr sycl::specialization_id same_name{13};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__spec_id_shim_[[SHIM_ID_2]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM_ID_2]]()>() {
 // CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -306,7 +306,7 @@ 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: static constexpr decltype(outer::same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return outer::same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -315,7 +315,7 @@ constexpr sycl::specialization_id same_name{14};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer9same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -327,7 +327,7 @@ constexpr sycl::specialization_id same_name{15};
 // 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -336,8 +336,8 @@ constexpr sycl::specialization_id same_name{15};
 // 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: static constexpr decltype(outer::__sycl_detail::__shim_[[SHIM_ID]]()) &__shim_[[SHIM_ID2:[0-9]+]]() {
+// CHECK-NEXT: return outer::__sycl_detail::__shim_[[SHIM_ID]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
 // CHECK-NEXT: } // namespace
@@ -345,7 +345,7 @@ constexpr sycl::specialization_id same_name{15};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID2]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -357,7 +357,7 @@ constexpr sycl::specialization_id same_name{16};
 // 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: static constexpr decltype(inner::same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return inner::same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -366,8 +366,8 @@ constexpr sycl::specialization_id same_name{16};
 // 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: static constexpr decltype(outer::__sycl_detail::__shim_[[SHIM_ID]]()) &__shim_[[SHIM_ID2:[0-9]+]]() {
+// CHECK-NEXT: return outer::__sycl_detail::__shim_[[SHIM_ID]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
 // CHECK-NEXT: } // namespace
@@ -375,7 +375,7 @@ constexpr sycl::specialization_id same_name{16};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID2]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID2]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_15outer12_GLOBAL__N_15inner9same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
diff --git a/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp b/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp
index 1299400e4d74d..378dcf4eeb2ec 100644
--- a/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp
+++ b/clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp
@@ -89,7 +89,7 @@ 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -98,7 +98,7 @@ constexpr sycl::specialization_id same_name{207};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -110,7 +110,7 @@ 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: static constexpr decltype(inner::same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return inner::same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -119,7 +119,7 @@ constexpr sycl::specialization_id same_name{208};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_15inner9same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -135,7 +135,7 @@ constexpr sycl::specialization_id same_name{209};
 // 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: static constexpr decltype(same_name) &__shim_[[SHIM_ID:[0-9]+]]() {
 // CHECK-NEXT: return same_name;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -146,7 +146,7 @@ constexpr sycl::specialization_id same_name{209};
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::__sycl_detail::__spec_id_shim_[[SHIM_ID]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::inner::__sycl_detail::__shim_[[SHIM_ID]]()>() {
 // CHECK-NEXT: return "____ZN5outer5inner12_GLOBAL__N_19same_nameE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
new file mode 100644
index 0000000000000..b6b4c8b4497ad
--- /dev/null
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -0,0 +1,109 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
+// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
+// RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
+
+#include "Inputs/sycl.hpp"
+
+using namespace cl::sycl::ext::oneapi;
+
+int main() {
+  cl::sycl::kernel_single_task<class first_kernel>([]() {});
+}
+
+// CHECK-HEADER: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-HEADER-NEXT: namespace sycl {
+// CHECK-HEADER-NEXT: namespace detail {
+// CHECK-HEADER-NEXT: namespace {
+// CHECK-HEADER-NEXT: class __sycl_device_global_registration {
+// CHECK-HEADER-NEXT: public:
+// CHECK-HEADER-NEXT:   __sycl_device_global_registration() noexcept;
+// CHECK-HEADER-NEXT: };
+// CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registerer;
+// CHECK-HEADER-NEXT: } // namespace
+// CHECK-HEADER: } // namespace detail
+// CHECK-HEADER: } // namespace sycl
+// CHECK-HEADER: } // __SYCL_INLINE_NAMESPACE(cl)
+
+// CHECK-FOOTER: #include <CL/sycl/detail/defines_elementary.hpp>
+// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
+
+// Shims go before the registration.
+// CHECK-FOOTER: namespace Foo {
+// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER-NEXT: namespace __sycl_detail {
+// CHECK-FOOTER-NEXT: static constexpr decltype(AnonNS) &__shim_[[SHIM0:[0-9]+]]() {
+// CHECK-FOOTER-NEXT:   return AnonNS;
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: } // namespace __sycl_detail
+// CHECK-FOOTER-NEXT: } // namespace 
+// CHECK-FOOTER-NEXT: } // namespace Foo
+// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER-NEXT: namespace __sycl_detail {
+// CHECK-FOOTER-NEXT: static constexpr decltype(HasVarTemplate::VarTempl<int>) &__shim_[[SHIM1:[0-9]+]]() {
+// CHECK-FOOTER-NEXT:   return HasVarTemplate::VarTempl<int>;
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: } // namespace __sycl_detail
+// CHECK-FOOTER-NEXT: } // namespace 
+
+
+// CHECK-FOOTER: namespace sycl::detail {
+// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+
+device_global<int> Basic;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Basic, "_Z5Basic");
+
+struct Wrapper {
+  static device_global<int> WrapperDevGlobal;
+};
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Wrapper::WrapperDevGlobal, "_ZN7Wrapper16WrapperDevGlobalE");
+
+template <typename T>
+struct WrapperTemplate {
+  static device_global<T> WrapperSpecID;
+};
+template class WrapperTemplate<int>;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::WrapperTemplate<int>::WrapperSpecID, "_ZN15WrapperTemplateIiE13WrapperSpecIDE");
+
+namespace Foo {
+device_global<int> NS;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::NS, "_ZN3Foo2NSE");
+
+inline namespace Bar {
+device_global<float> InlineNS;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::InlineNS, "_ZN3Foo3Bar8InlineNSE");
+
+device_global<int> NS;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::Bar::NS, "_ZN3Foo3Bar2NSE");
+
+struct Wrapper {
+  static device_global<int> WrapperDevGlobal;
+};
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::Wrapper::WrapperDevGlobal, "_ZN3Foo3Bar7Wrapper16WrapperDevGlobalE");
+
+template <typename T>
+struct WrapperTemplate {
+  static device_global<T> WrapperDevGlobal;
+};
+template class WrapperTemplate<float>;
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::WrapperTemplate<float>::WrapperDevGlobal, "_ZN3Foo3Bar15WrapperTemplateIfE16WrapperDevGlobalE");
+} // namespace Bar
+
+namespace {
+device_global<int> AnonNS;
+}
+// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::__sycl_detail::__shim_[[SHIM0]](), "____ZN3Foo12_GLOBAL__N_16AnonNSE");
+
+} // namespace Foo
+
+// Validate that variable templates work correctly.
+namespace {
+struct HasVarTemplate {
+  constexpr HasVarTemplate(){}
+  template<typename T>
+  static constexpr device_global<T> VarTempl{};
+};
+
+}
+const auto x = HasVarTemplate::VarTempl<int>.get();
+// CHECK-FOOTER-NEXT: device_global_map::add(&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");
diff --git a/clang/test/CodeGenSYCL/integration_footer.cpp b/clang/test/CodeGenSYCL/integration_footer.cpp
index 387399588556a..a9fd6dcb6982e 100644
--- a/clang/test/CodeGenSYCL/integration_footer.cpp
+++ b/clang/test/CodeGenSYCL/integration_footer.cpp
@@ -127,7 +127,7 @@ specialization_id<int> AnonNSSpecID;
 // CHECK: namespace Foo {
 // CHECK: namespace {
 // CHECK-NEXT: namespace __sycl_detail {
-// CHECK-NEXT: static constexpr decltype(AnonNSSpecID) &__spec_id_shim_[[SHIM0:[0-9]+]]() {
+// CHECK-NEXT: static constexpr decltype(AnonNSSpecID) &__shim_[[SHIM0:[0-9]+]]() {
 // CHECK-NEXT: return AnonNSSpecID;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -137,7 +137,7 @@ specialization_id<int> AnonNSSpecID;
 // CHECK: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::Foo::__sycl_detail::__spec_id_shim_[[SHIM0]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::Foo::__sycl_detail::__shim_[[SHIM0]]()>() {
 // CHECK-NEXT: return "____ZN3Foo12_GLOBAL__N_112AnonNSSpecIDE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail
@@ -174,7 +174,7 @@ struct HasVarTemplate {
 auto x = HasVarTemplate::VarTempl<int, 2>.getDefaultValue();
 // CHECK: namespace {
 // CHECK-NEXT: namespace __sycl_detail
-// CHECK-NEXT: static constexpr decltype(HasVarTemplate::VarTempl<int, 2>) &__spec_id_shim_[[SHIM1:[0-9]+]]() {
+// CHECK-NEXT: static constexpr decltype(HasVarTemplate::VarTempl<int, 2>) &__shim_[[SHIM1:[0-9]+]]() {
 // CHECK-NEXT: return HasVarTemplate::VarTempl<int, 2>;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
@@ -183,7 +183,7 @@ auto x = HasVarTemplate::VarTempl<int, 2>.getDefaultValue();
 // CHECK-NEXT: namespace sycl {
 // CHECK-NEXT: namespace detail {
 // CHECK-NEXT: template<>
-// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__spec_id_shim_[[SHIM1]]()>() {
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM1]]()>() {
 // CHECK-NEXT: return "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiLi2EEE";
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace detail

From 28c468607a052f824e4383eba5b9306d456b7aa7 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 13:20:28 +0300
Subject: [PATCH 02/11] Extend the test

---
 clang/lib/Sema/SemaSYCL.cpp                   |  2 +-
 .../CL/sycl/detail/defines_elementary.hpp     |  8 +++
 .../CL/sycl/detail/device_global_map.hpp      | 13 +++++
 .../Inputs/CL/sycl/detail/kernel_desc.hpp     | 49 +++++++++++++++++++
 .../device_global_int_footer_header.cpp       | 33 +++++++------
 ...a-attribute-supported-attributes-list.test |  1 +
 6 files changed, 91 insertions(+), 15 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
 create mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
 create mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 92d5901c49b71..c06a4166ec5b0 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -5101,7 +5101,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     }
     VarRefNameOS.flush();
     if (Util::isSyclDeviceGlobalType(VD->getType())) {
-      DeviceGlobOS << "&";
+      DeviceGlobOS << "(void *)&";
       DeviceGlobOS << VarRefName;
     } else {
       OS << VarRefName;
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
new file mode 100644
index 0000000000000..00322e60ffc51
--- /dev/null
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
@@ -0,0 +1,8 @@
+#pragma once
+
+#ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
+#define __SYCL_INLINE_NAMESPACE(X) inline namespace X
+#else
+#define __SYCL_INLINE_NAMESPACE(X) namespace X
+#endif // __SYCL_DISABLE_NAMESPACE_INLINE__
+#define __SYCL_DLL_LOCAL
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
new file mode 100644
index 0000000000000..17d690daeef54
--- /dev/null
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
@@ -0,0 +1,13 @@
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+namespace device_global_map {
+
+void add(void *DeviceGlobalPtr, const char *UniqueId);
+
+} // namespace device_global_map
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
new file mode 100644
index 0000000000000..c1812a0324868
--- /dev/null
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
@@ -0,0 +1,49 @@
+#pragma once
+
+#include <CL/sycl/detail/defines_elementary.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+  namespace sycl {
+  namespace detail {
+
+#ifndef __SYCL_DEVICE_ONLY__
+#define _Bool bool
+#endif
+
+  // kernel parameter kinds
+  enum class kernel_param_kind_t {
+    kind_accessor = 0,
+    kind_std_layout = 1, // standard layout object parameters
+    kind_sampler = 2,
+    kind_pointer = 3,
+    kind_specialization_constants_buffer = 4,
+    kind_stream = 5,
+    kind_invalid = 0xf, // not a valid kernel kind
+  };
+
+  // describes a kernel parameter
+  struct kernel_param_desc_t {
+    // parameter kind
+    kernel_param_kind_t kind;
+    // kind == kind_std_layout
+    //   parameter size in bytes (includes padding for structs)
+    // kind == kind_accessor
+    //   access target; possible access targets are defined in access/access.hpp
+    int info;
+    // offset of the captured value of the parameter in the lambda or function
+    // object
+    int offset;
+  };
+
+  template <class KernelNameType> struct KernelInfo {
+    static constexpr unsigned getNumParams() { return 0; }
+    static const kernel_param_desc_t &getParamDesc(int) {
+      static kernel_param_desc_t Dummy;
+      return Dummy;
+    }
+    static constexpr const char *getName() { return ""; }
+    static constexpr bool isESIMD() { return 0; }
+  };
+  } // namespace detail
+  } // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index b6b4c8b4497ad..4d3a33874532e 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -1,8 +1,16 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
 // RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
 
-#include "Inputs/sycl.hpp"
+// Try and compile all this stuff.
+// RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
+
+
+// This test checks that integration header and footer are emitted correctly
+// for device_global variables. It also checks that emitted costructs
+// are syntactically correct.
+
+#include "sycl.hpp"
 
 using namespace cl::sycl::ext::oneapi;
 
@@ -51,48 +59,45 @@ int main() {
 // CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
 
 device_global<int> Basic;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Basic, "_Z5Basic");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Basic, "_Z5Basic");
 
 struct Wrapper {
   static device_global<int> WrapperDevGlobal;
 };
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Wrapper::WrapperDevGlobal, "_ZN7Wrapper16WrapperDevGlobalE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Wrapper::WrapperDevGlobal, "_ZN7Wrapper16WrapperDevGlobalE");
 
 template <typename T>
 struct WrapperTemplate {
   static device_global<T> WrapperSpecID;
 };
 template class WrapperTemplate<int>;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::WrapperTemplate<int>::WrapperSpecID, "_ZN15WrapperTemplateIiE13WrapperSpecIDE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::WrapperTemplate<int>::WrapperSpecID, "_ZN15WrapperTemplateIiE13WrapperSpecIDE");
 
 namespace Foo {
 device_global<int> NS;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::NS, "_ZN3Foo2NSE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::NS, "_ZN3Foo2NSE");
 
 inline namespace Bar {
 device_global<float> InlineNS;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::InlineNS, "_ZN3Foo3Bar8InlineNSE");
-
-device_global<int> NS;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::Bar::NS, "_ZN3Foo3Bar2NSE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::InlineNS, "_ZN3Foo3Bar8InlineNSE");
 
 struct Wrapper {
   static device_global<int> WrapperDevGlobal;
 };
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::Wrapper::WrapperDevGlobal, "_ZN3Foo3Bar7Wrapper16WrapperDevGlobalE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::Wrapper::WrapperDevGlobal, "_ZN3Foo3Bar7Wrapper16WrapperDevGlobalE");
 
 template <typename T>
 struct WrapperTemplate {
   static device_global<T> WrapperDevGlobal;
 };
 template class WrapperTemplate<float>;
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::WrapperTemplate<float>::WrapperDevGlobal, "_ZN3Foo3Bar15WrapperTemplateIfE16WrapperDevGlobalE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::WrapperTemplate<float>::WrapperDevGlobal, "_ZN3Foo3Bar15WrapperTemplateIfE16WrapperDevGlobalE");
 } // namespace Bar
 
 namespace {
 device_global<int> AnonNS;
 }
-// CHECK-FOOTER-NEXT: device_global_map::add(&::Foo::__sycl_detail::__shim_[[SHIM0]](), "____ZN3Foo12_GLOBAL__N_16AnonNSE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::__sycl_detail::__shim_[[SHIM0]](), "____ZN3Foo12_GLOBAL__N_16AnonNSE");
 
 } // namespace Foo
 
@@ -106,4 +111,4 @@ struct HasVarTemplate {
 
 }
 const auto x = HasVarTemplate::VarTempl<int>.get();
-// CHECK-FOOTER-NEXT: device_global_map::add(&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 065375f39483e..4c1c3e14bca59 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -157,6 +157,7 @@
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLDeviceGlobal (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function)

From 2942920268580ada19e7d15a3dedd92bea7e4c43 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 13:40:04 +0300
Subject: [PATCH 03/11] Make clang-format happy

---
 clang/lib/Sema/SemaSYCL.cpp                        |  9 +++++----
 .../Inputs/CL/sycl/detail/device_global_map.hpp    | 14 +++++++-------
 clang/test/CodeGenSYCL/Inputs/sycl.hpp             |  3 ++-
 .../device_global_int_footer_header.cpp            | 12 +++++-------
 4 files changed, 19 insertions(+), 19 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c06a4166ec5b0..2595615453678 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4979,8 +4979,8 @@ static std::string EmitSpecIdShim(raw_ostream &OS, unsigned &ShimCounter,
   // Print opening-namespace
   PrintNamespaces(OS, Decl::castToDeclContext(AnonNS));
   OS << "namespace __sycl_detail {\n";
-  OS << "static constexpr decltype(" << LastShim << ") &__shim_"
-     << ShimCounter << "() {\n";
+  OS << "static constexpr decltype(" << LastShim << ") &__shim_" << ShimCounter
+     << "() {\n";
   OS << "  return " << LastShim << ";\n";
   OS << "}\n";
   OS << "} // namespace __sycl_detail\n";
@@ -5109,7 +5109,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
     if (Util::isSyclDeviceGlobalType(VD->getType())) {
       DeviceGlobOS << ", \"";
-      DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD);
+      DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
+                                                          VD);
       DeviceGlobOS << "\");\n";
     } else {
       OS << ">() {\n";
@@ -5188,7 +5189,7 @@ bool Util::isSyclDeviceGlobalType(QualType Ty) {
   if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecTy)) {
     ClassTemplateDecl *Template = CTSD->getSpecializedTemplate();
     if (CXXRecordDecl *RD = Template->getTemplatedDecl())
-        return RD->hasAttr<SYCLDeviceGlobalAttr>();
+      return RD->hasAttr<SYCLDeviceGlobalAttr>();
   }
   return RecTy->hasAttr<SYCLDeviceGlobalAttr>();
 }
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
index 17d690daeef54..a2897f48b419a 100644
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
@@ -1,13 +1,13 @@
 #pragma once
 
 __SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl {
-namespace detail {
-namespace device_global_map {
+  namespace sycl {
+  namespace detail {
+  namespace device_global_map {
 
-void add(void *DeviceGlobalPtr, const char *UniqueId);
+  void add(void *DeviceGlobalPtr, const char *UniqueId);
 
-} // namespace device_global_map
-} // namespace detail
-} // namespace sycl
+  } // namespace device_global_map
+  } // namespace detail
+  } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 465d97b20bcce..143e8743a4822 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -133,8 +133,9 @@ struct no_alias {
 template <typename T>
 class [[__sycl_detail__::device_global]] device_global {
 public:
-  const T& get() const noexcept { return *Data; }
+  const T &get() const noexcept { return *Data; }
   device_global() = default;
+
 private:
   T *Data;
 };
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index 4d3a33874532e..10da802eee8a5 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -5,7 +5,6 @@
 // Try and compile all this stuff.
 // RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
 
-
 // This test checks that integration header and footer are emitted correctly
 // for device_global variables. It also checks that emitted costructs
 // are syntactically correct.
@@ -43,7 +42,7 @@ int main() {
 // CHECK-FOOTER-NEXT:   return AnonNS;
 // CHECK-FOOTER-NEXT: }
 // CHECK-FOOTER-NEXT: } // namespace __sycl_detail
-// CHECK-FOOTER-NEXT: } // namespace 
+// CHECK-FOOTER-NEXT: } // namespace
 // CHECK-FOOTER-NEXT: } // namespace Foo
 // CHECK-FOOTER-NEXT: namespace {
 // CHECK-FOOTER-NEXT: namespace __sycl_detail {
@@ -51,8 +50,7 @@ int main() {
 // CHECK-FOOTER-NEXT:   return HasVarTemplate::VarTempl<int>;
 // CHECK-FOOTER-NEXT: }
 // CHECK-FOOTER-NEXT: } // namespace __sycl_detail
-// CHECK-FOOTER-NEXT: } // namespace 
-
+// CHECK-FOOTER-NEXT: } // namespace
 
 // CHECK-FOOTER: namespace sycl::detail {
 // CHECK-FOOTER-NEXT: namespace {
@@ -104,11 +102,11 @@ device_global<int> AnonNS;
 // Validate that variable templates work correctly.
 namespace {
 struct HasVarTemplate {
-  constexpr HasVarTemplate(){}
-  template<typename T>
+  constexpr HasVarTemplate() {}
+  template <typename T>
   static constexpr device_global<T> VarTempl{};
 };
 
-}
+} // namespace
 const auto x = HasVarTemplate::VarTempl<int>.get();
 // CHECK-FOOTER-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");

From 25383961fdcd04b3a9ea99930dd31be00752effe Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 14:06:42 +0300
Subject: [PATCH 04/11] Still not happy enough

---
 clang/lib/Sema/SemaSYCL.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2595615453678..717656cf8736a 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -5073,7 +5073,6 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     if (!EmittedFirstSpecConstant && !DeviceGlobalsEmitted)
       OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
 
-
     Visited.insert(VD);
     std::string TopShim = EmitSpecIdShims(OS, ShimCounter, Policy, VD);
     if (Util::isSyclDeviceGlobalType(VD->getType())) {

From 837f7817e1c26887d0b1a52d4f435ad14b89bde4 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 17:28:24 +0300
Subject: [PATCH 05/11] Add test checking that spec ids work with device
 globals

---
 clang/lib/Sema/SemaSYCL.cpp                   |   4 +-
 .../Inputs/CL/sycl/detail/kernel_desc.hpp     |   5 +
 .../CL/sycl/detail/spec_const_integration.hpp |  15 +
 .../device_global_int_footer_header.cpp       |   2 +-
 .../device_globals_with_spec_ids.cpp          | 267 ++++++++++++++++++
 5 files changed, 289 insertions(+), 4 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
 create mode 100644 clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 717656cf8736a..99dd1ccd9c97d 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -5076,9 +5076,6 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     Visited.insert(VD);
     std::string TopShim = EmitSpecIdShims(OS, ShimCounter, Policy, VD);
     if (Util::isSyclDeviceGlobalType(VD->getType())) {
-      if (!DeviceGlobalsEmitted)
-        OS << "#include <CL/sycl/detail/device_global_map.hpp>\n";
-
       DeviceGlobalsEmitted = true;
       DeviceGlobOS << "device_global_map::add(";
     } else {
@@ -5127,6 +5124,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
 
   if (DeviceGlobalsEmitted) {
+    OS << "#include <CL/sycl/detail/device_global_map.hpp>\n";
     DeviceGlobOS.flush();
     OS << "namespace sycl::detail {\n";
     OS << "namespace {\n";
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
index c1812a0324868..c4d9b998e3cb2 100644
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
@@ -6,6 +6,11 @@ __SYCL_INLINE_NAMESPACE(cl) {
   namespace sycl {
   namespace detail {
 
+#if __cplusplus >= 201703L
+  template <auto &SpecName> const char *get_spec_constant_symbolic_ID_impl();
+  template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
+#endif
+
 #ifndef __SYCL_DEVICE_ONLY__
 #define _Bool bool
 #endif
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
new file mode 100644
index 0000000000000..7a97db07f4036
--- /dev/null
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
@@ -0,0 +1,15 @@
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+#if __cplusplus >= 201703L
+// Translates SYCL 2020 `specialization_id` to a unique symbolic identifier
+// which is used internally by the toolchain
+template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
+  return get_spec_constant_symbolic_ID_impl<SpecName>();
+}
+#endif
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index 10da802eee8a5..47a5445ef3860 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -32,7 +32,6 @@ int main() {
 // CHECK-HEADER: } // __SYCL_INLINE_NAMESPACE(cl)
 
 // CHECK-FOOTER: #include <CL/sycl/detail/defines_elementary.hpp>
-// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
 
 // Shims go before the registration.
 // CHECK-FOOTER: namespace Foo {
@@ -52,6 +51,7 @@ int main() {
 // CHECK-FOOTER-NEXT: } // namespace __sycl_detail
 // CHECK-FOOTER-NEXT: } // namespace
 
+// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
 // CHECK-FOOTER: namespace sycl::detail {
 // CHECK-FOOTER-NEXT: namespace {
 // CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
diff --git a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
new file mode 100644
index 0000000000000..e50531691db5a
--- /dev/null
+++ b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
@@ -0,0 +1,267 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
+// RUN: FileCheck -input-file=%t.footer.h %s
+
+// Try and compile all this stuff.
+// RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
+
+// This test checks that integration footer is emitted correctly if both
+// spec constants and device globals are used.
+
+#include "sycl.hpp"
+
+using namespace cl;
+int main() {
+  cl::sycl::kernel_single_task<class first_kernel>([]() {});
+}
+
+// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
+constexpr sycl::specialization_id a{2};
+// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::a>() {
+// CHECK-NEXT: return "____ZL1a";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+sycl::ext::oneapi::device_global<int> b;
+
+struct Wrapper {
+  static constexpr sycl::specialization_id a{18};
+  static sycl::ext::oneapi::device_global<float> b;
+};
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::Wrapper::a>() {
+// CHECK-NEXT:   return "_ZN7Wrapper1aE";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+
+template <typename T>
+struct TemplateWrapper {
+  static constexpr sycl::specialization_id<T> a{18};
+  static sycl::ext::oneapi::device_global<T> b;
+};
+
+template class TemplateWrapper<float>;
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::TemplateWrapper<float>::a>() {
+// CHECK-NEXT:   return "_ZN15TemplateWrapperIfE1aE";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+
+namespace {
+constexpr sycl::specialization_id a{2};
+sycl::ext::oneapi::device_global<int> b;
+} // namespace
+
+// CHECK: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(a) &__shim_[[SHIM0:[0-9]+]]() {
+// CHECK-NEXT:   return a;
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+
+// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM0]]()>() {
+// CHECK-NEXT:   return "____ZN12_GLOBAL__N_11aE";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+
+// CHECK: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(b) &__shim_[[SHIM1:[0-9]+]]() {
+// CHECK-NEXT:   return b;
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+
+namespace outer {
+namespace {
+namespace inner {
+namespace {
+constexpr sycl::specialization_id a{2};
+// CHECK: namespace outer {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace inner {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(a) &__shim_[[SHIM2:[0-9]+]]() {
+// CHECK-NEXT:   return a;
+// 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::__shim_[[SHIM2]]()) &__shim_[[SHIM3:[0-9]+]]() {
+// CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM2]]();
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace outer
+// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM3]]()>() {
+// CHECK-NEXT:   return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_11aE";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+sycl::ext::oneapi::device_global<int> b;
+// CHECK: namespace outer {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace inner {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(b) &__shim_[[SHIM4:[0-9]+]]() {
+// CHECK-NEXT:   return b;
+// 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::__shim_[[SHIM4]]()) &__shim_[[SHIM5:[0-9]+]]() {
+// CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM4]]();
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace outer
+struct Wrapper {
+  static constexpr sycl::specialization_id a{18};
+  static sycl::ext::oneapi::device_global<int> b;
+  static sycl::ext::oneapi::device_global<float> c;
+};
+// CHECK: namespace outer {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace inner {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(Wrapper::a) &__shim_[[SHIM6:[0-9]+]]() {
+// CHECK-NEXT:   return Wrapper::a;
+// 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::__shim_[[SHIM6]]()) &__shim_[[SHIM7:[0-9]+]]() {
+// CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM6]]();
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace outer
+// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: template<>
+// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM7]]()>() {
+// CHECK-NEXT:   return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1aE";
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace detail
+// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
+
+// CHECK-NEXT: namespace outer {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace inner {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(Wrapper::b) &__shim_[[SHIM8:[0-9]+]]() {
+// CHECK-NEXT:   return Wrapper::b;
+// 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::__shim_[[SHIM8]]()) &__shim_[[SHIM9:[0-9]+]]() {
+// CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM8]]();
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace outer
+// CHECK-NEXT: namespace outer {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace inner {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: namespace __sycl_detail {
+// CHECK-NEXT: static constexpr decltype(Wrapper::c) &__shim_[[SHIM10:[0-9]+]]() {
+// CHECK-NEXT:   return Wrapper::c;
+// 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::__shim_[[SHIM10]]()) &__shim_[[SHIM11:[0-9]+]]() {
+// CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM10]]();
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace __sycl_detail
+// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace outer
+
+// FIXME: Shims don't work with templated wrapper classes for some reason
+// template <typename T>
+// struct TemplateWrapper {
+//   static constexpr sycl::specialization_id<T> a{18};
+//   static sycl::ext::oneapi::device_global<T> b;
+// };
+// 
+// template class TemplateWrapper<float>;
+
+}
+}
+}
+}
+
+// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>
+// CHECK-NEXT: #include <CL/sycl/detail/device_global_map.hpp>
+// CHECK-NEXT: namespace sycl::detail {
+// CHECK-NEXT: namespace {
+// CHECK-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+// CHECK-NEXT: device_global_map::add((void *)&::b, "_Z1b");
+// CHECK-NEXT: device_global_map::add((void *)&::Wrapper::b, "_ZN7Wrapper1bE");
+// CHECK-NEXT: device_global_map::add((void *)&::TemplateWrapper<float>::b, "_ZN15TemplateWrapperIfE1bE");
+// CHECK-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_11bE");
+// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM5]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_11bE");
+// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM9]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1bE");
+// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM11]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1cE");
+// CHECK-NEXT: }
+// CHECK-NEXT: } // namespace (unnamed)
+// CHECK-NEXT: } // namespace sycl::detail

From c692020d78c5a4121c7abe234aeb57cce709851c Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 15 Feb 2022 17:46:30 +0300
Subject: [PATCH 06/11] Make clang-format happy again

---
 .../CL/sycl/detail/spec_const_integration.hpp | 18 ++++----
 .../device_globals_with_spec_ids.cpp          | 44 +++++++++----------
 2 files changed, 31 insertions(+), 31 deletions(-)

diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
index 7a97db07f4036..c3d865352ffa5 100644
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
@@ -1,15 +1,15 @@
 __SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl {
-namespace detail {
+  namespace sycl {
+  namespace detail {
 
 #if __cplusplus >= 201703L
-// Translates SYCL 2020 `specialization_id` to a unique symbolic identifier
-// which is used internally by the toolchain
-template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
-  return get_spec_constant_symbolic_ID_impl<SpecName>();
-}
+  // Translates SYCL 2020 `specialization_id` to a unique symbolic identifier
+  // which is used internally by the toolchain
+  template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
+    return get_spec_constant_symbolic_ID_impl<SpecName>();
+  }
 #endif
 
-} // namespace detail
-} // namespace sycl
+  } // namespace detail
+  } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
index e50531691db5a..bef9f71998d11 100644
--- a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
+++ b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
@@ -72,7 +72,7 @@ sycl::ext::oneapi::device_global<int> b;
 // CHECK-NEXT:   return a;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 
 // CHECK: __SYCL_INLINE_NAMESPACE(cl) {
 // CHECK-NEXT: namespace sycl {
@@ -91,7 +91,7 @@ sycl::ext::oneapi::device_global<int> b;
 // CHECK-NEXT:   return b;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 
 namespace outer {
 namespace {
@@ -107,9 +107,9 @@ constexpr sycl::specialization_id a{2};
 // CHECK-NEXT:   return a;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace inner
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -118,7 +118,7 @@ constexpr sycl::specialization_id a{2};
 // CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM2]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
 // CHECK-NEXT: namespace sycl {
@@ -140,9 +140,9 @@ sycl::ext::oneapi::device_global<int> b;
 // CHECK-NEXT:   return b;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace inner
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -151,7 +151,7 @@ sycl::ext::oneapi::device_global<int> b;
 // CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM4]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 struct Wrapper {
   static constexpr sycl::specialization_id a{18};
@@ -167,9 +167,9 @@ struct Wrapper {
 // CHECK-NEXT:   return Wrapper::a;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace inner
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -178,7 +178,7 @@ struct Wrapper {
 // CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM6]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
 // CHECK-NEXT: namespace sycl {
@@ -200,9 +200,9 @@ struct Wrapper {
 // CHECK-NEXT:   return Wrapper::b;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace inner
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -211,7 +211,7 @@ struct Wrapper {
 // CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM8]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -222,9 +222,9 @@ struct Wrapper {
 // CHECK-NEXT:   return Wrapper::c;
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace inner
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 // CHECK-NEXT: namespace outer {
 // CHECK-NEXT: namespace {
@@ -233,7 +233,7 @@ struct Wrapper {
 // CHECK-NEXT:   return inner::__sycl_detail::__shim_[[SHIM10]]();
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace __sycl_detail
-// CHECK-NEXT: } // namespace 
+// CHECK-NEXT: } // namespace
 // CHECK-NEXT: } // namespace outer
 
 // FIXME: Shims don't work with templated wrapper classes for some reason
@@ -242,13 +242,13 @@ struct Wrapper {
 //   static constexpr sycl::specialization_id<T> a{18};
 //   static sycl::ext::oneapi::device_global<T> b;
 // };
-// 
+//
 // template class TemplateWrapper<float>;
 
-}
-}
-}
-}
+} // namespace
+} // namespace inner
+} // namespace
+} // namespace outer
 
 // CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>
 // CHECK-NEXT: #include <CL/sycl/detail/device_global_map.hpp>

From 0f288d7e6ff15469283f0307af4739ac838bb1b3 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Mon, 21 Feb 2022 11:34:35 +0300
Subject: [PATCH 07/11] Fix copy-paste error and add the prefix

---
 .../CodeGenSYCL/device_global_int_footer_header.cpp    | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index 47a5445ef3860..bb9ce2cd44962 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h -fsycl-unique-prefix=THE_PREFIX %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
 // RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
 
@@ -66,10 +66,10 @@ struct Wrapper {
 
 template <typename T>
 struct WrapperTemplate {
-  static device_global<T> WrapperSpecID;
+  static device_global<T> WrapperDevGlobal;
 };
 template class WrapperTemplate<int>;
-// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::WrapperTemplate<int>::WrapperSpecID, "_ZN15WrapperTemplateIiE13WrapperSpecIDE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::WrapperTemplate<int>::WrapperDevGlobal, "_ZN15WrapperTemplateIiE16WrapperDevGlobalE");
 
 namespace Foo {
 device_global<int> NS;
@@ -95,7 +95,7 @@ template class WrapperTemplate<float>;
 namespace {
 device_global<int> AnonNS;
 }
-// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::__sycl_detail::__shim_[[SHIM0]](), "____ZN3Foo12_GLOBAL__N_16AnonNSE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Foo::__sycl_detail::__shim_[[SHIM0]](), "THE_PREFIX____ZN3Foo12_GLOBAL__N_16AnonNSE");
 
 } // namespace Foo
 
@@ -109,4 +109,4 @@ struct HasVarTemplate {
 
 } // namespace
 const auto x = HasVarTemplate::VarTempl<int>.get();
-// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "THE_PREFIX____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");

From 683d7eb67e9c985e6139dbfb9862d44555040729 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Thu, 24 Feb 2022 14:24:07 +0300
Subject: [PATCH 08/11] Apply some review feedback

---
 clang/include/clang/Sema/Sema.h                |  2 +-
 clang/lib/Sema/Sema.cpp                        |  4 ++--
 clang/lib/Sema/SemaSYCL.cpp                    | 18 +++++++++---------
 .../device_global_int_footer_header.cpp        |  6 +++---
 .../device_globals_with_spec_ids.cpp           |  2 +-
 5 files changed, 16 insertions(+), 16 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index d3091268b2634..69093174864d5 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -440,7 +440,7 @@ class SYCLIntegrationFooter {
   SYCLIntegrationFooter(Sema &S) : S(S) {}
   bool emit(StringRef MainSrc);
   void addVarDecl(const VarDecl *VD);
-  bool metSYCLDeviceGlobals() { return DeviceGlobalsEmitted; }
+  bool isDeviceGlobalsEmitted() { return DeviceGlobalsEmitted; }
 
 private:
   bool emit(raw_ostream &O);
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 115290145d8a4..17ccbe32a93f6 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1101,8 +1101,8 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
     // Set the names of the kernels, now that the names have settled down. This
     // needs to happen before we generate the integration headers.
     SetSYCLKernelNames();
-    // Make sure that the footer emitted before header, since only after the
-    // footer is emitted it is known that translation unit contains device
+    // Make sure that the footer is emitted before header, since only after the
+    // footer is emitted is it known that translation unit contains device
     // global variables.
     if (SyclIntFooter != nullptr)
       SyclIntFooter->emit(getLangOpts().SYCLIntFooter);
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 99dd1ccd9c97d..52152fabdc6f9 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4684,14 +4684,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   // whose sole purpose is to run its constructor before the application's
   // main() function.
 
-  if (S.getSyclIntegrationFooter().metSYCLDeviceGlobals()) {
+  if (S.getSyclIntegrationFooter().isDeviceGlobalsEmitted()) {
     O << "namespace {\n";
 
     O << "class __sycl_device_global_registration {\n";
     O << "public:\n";
     O << "  __sycl_device_global_registration() noexcept;\n";
     O << "};\n";
-    O << "__sycl_device_global_registration __sycl_device_global_registerer;\n";
+    O << "__sycl_device_global_registration __sycl_device_global_registrar;\n";
 
     O << "} // namespace\n";
 
@@ -4971,7 +4971,7 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
       [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC);
 }
 
-static std::string EmitSpecIdShim(raw_ostream &OS, unsigned &ShimCounter,
+static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter,
                                   const std::string &LastShim,
                                   const NamespaceDecl *AnonNS) {
   std::string NewShimName =
@@ -4991,7 +4991,7 @@ static std::string EmitSpecIdShim(raw_ostream &OS, unsigned &ShimCounter,
 }
 
 // Emit the list of shims required for a DeclContext, calls itself recursively.
-static void EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter,
+static void EmitShims(raw_ostream &OS, unsigned &ShimCounter,
                             const DeclContext *DC,
                             std::string &NameForLastShim) {
   if (DC->isTranslationUnit()) {
@@ -5007,7 +5007,7 @@ static void EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter,
   } else if (const auto *ND = dyn_cast<NamespaceDecl>(CurDecl)) {
     if (ND->isAnonymousNamespace()) {
       // Print current shim, reset 'name for last shim'.
-      NameForLastShim = EmitSpecIdShim(OS, ShimCounter, NameForLastShim, ND);
+      NameForLastShim = EmitShim(OS, ShimCounter, NameForLastShim, ND);
     } else {
       NameForLastShim = ND->getNameAsString() + "::" + NameForLastShim;
     }
@@ -5021,13 +5021,13 @@ static void EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter,
            "Unhandled decl type");
   }
 
-  EmitSpecIdShims(OS, ShimCounter, CurDecl->getDeclContext(), NameForLastShim);
+  EmitShims(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,
+static std::string EmitShims(raw_ostream &OS, unsigned &ShimCounter,
                                    PrintingPolicy &Policy, const VarDecl *VD) {
   if (!VD->isInAnonymousNamespace())
     return "";
@@ -5036,7 +5036,7 @@ static std::string EmitSpecIdShims(raw_ostream &OS, unsigned &ShimCounter,
   VD->getNameForDiagnostic(stream, Policy, false);
   stream.flush();
 
-  EmitSpecIdShims(OS, ShimCounter, VD->getDeclContext(), RelativeName);
+  EmitShims(OS, ShimCounter, VD->getDeclContext(), RelativeName);
   return RelativeName;
 }
 
@@ -5074,7 +5074,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
 
     Visited.insert(VD);
-    std::string TopShim = EmitSpecIdShims(OS, ShimCounter, Policy, VD);
+    std::string TopShim = EmitShims(OS, ShimCounter, Policy, VD);
     if (Util::isSyclDeviceGlobalType(VD->getType())) {
       DeviceGlobalsEmitted = true;
       DeviceGlobOS << "device_global_map::add(";
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index bb9ce2cd44962..44ffe0f72f281 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -2,11 +2,11 @@
 // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
 // RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
 
-// Try and compile all this stuff.
+// Try and compile generated integration header and footer on host.
 // RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
 
 // This test checks that integration header and footer are emitted correctly
-// for device_global variables. It also checks that emitted costructs
+// for device_global variables. It also checks that emitted constructs
 // are syntactically correct.
 
 #include "sycl.hpp"
@@ -25,7 +25,7 @@ int main() {
 // CHECK-HEADER-NEXT: public:
 // CHECK-HEADER-NEXT:   __sycl_device_global_registration() noexcept;
 // CHECK-HEADER-NEXT: };
-// CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registerer;
+// CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registrar;
 // CHECK-HEADER-NEXT: } // namespace
 // CHECK-HEADER: } // namespace detail
 // CHECK-HEADER: } // namespace sycl
diff --git a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
index bef9f71998d11..674401c0d79a6 100644
--- a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
+++ b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.footer.h %s
 
-// Try and compile all this stuff.
+// Try and compile generated integration header and footer on host.
 // RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
 
 // This test checks that integration footer is emitted correctly if both

From 3c3b5e54ac257e0b1753e2de17a65bd3d18efa29 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Thu, 24 Feb 2022 15:58:26 +0300
Subject: [PATCH 09/11] Arrange code in SemaSYCL

---
 clang/lib/Sema/SemaSYCL.cpp | 50 ++++++++++++++++---------------------
 1 file changed, 22 insertions(+), 28 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 52152fabdc6f9..dbc9a2586d2f9 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4972,8 +4972,8 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
 }
 
 static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter,
-                                  const std::string &LastShim,
-                                  const NamespaceDecl *AnonNS) {
+                            const std::string &LastShim,
+                            const NamespaceDecl *AnonNS) {
   std::string NewShimName =
       "__sycl_detail::__shim_" + std::to_string(ShimCounter) + "()";
   // Print opening-namespace
@@ -4992,8 +4992,7 @@ static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter,
 
 // Emit the list of shims required for a DeclContext, calls itself recursively.
 static void EmitShims(raw_ostream &OS, unsigned &ShimCounter,
-                            const DeclContext *DC,
-                            std::string &NameForLastShim) {
+                      const DeclContext *DC, std::string &NameForLastShim) {
   if (DC->isTranslationUnit()) {
     NameForLastShim = "::" + NameForLastShim;
     return;
@@ -5028,7 +5027,7 @@ static void EmitShims(raw_ostream &OS, unsigned &ShimCounter,
 // Returns a string containing the FQN of the 'top most' shim, including its
 // function call parameters.
 static std::string EmitShims(raw_ostream &OS, unsigned &ShimCounter,
-                                   PrintingPolicy &Policy, const VarDecl *VD) {
+                             PrintingPolicy &Policy, const VarDecl *VD) {
   if (!VD->isInAnonymousNamespace())
     return "";
   std::string RelativeName;
@@ -5068,7 +5067,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     if (llvm::find(Visited, VD) != Visited.end())
       continue;
 
-    // We only want to emit the #includes if we have a spec-constant that needs
+    // We only want to emit the #includes if we have a variable that needs
     // them, so emit this one on the first time through the loop.
     if (!EmittedFirstSpecConstant && !DeviceGlobalsEmitted)
       OS << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
@@ -5078,6 +5077,17 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     if (Util::isSyclDeviceGlobalType(VD->getType())) {
       DeviceGlobalsEmitted = true;
       DeviceGlobOS << "device_global_map::add(";
+      DeviceGlobOS << "(void *)&";
+      if (VD->isInAnonymousNamespace()) {
+        DeviceGlobOS << TopShim;
+      } else {
+        DeviceGlobOS << "::";
+        VD->getNameForDiagnostic(DeviceGlobOS, Policy, true);
+      }
+      DeviceGlobOS << ", \"";
+      DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
+                                                          VD);
+      DeviceGlobOS << "\");\n";
     } else {
       EmittedFirstSpecConstant = true;
       OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
@@ -5085,30 +5095,14 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       OS << "namespace detail {\n";
       OS << "template<>\n";
       OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
-    }
 
-    std::string VarRefName;
-    llvm::raw_string_ostream VarRefNameOS(VarRefName);
-    if (VD->isInAnonymousNamespace()) {
-      VarRefNameOS << TopShim;
-    } else {
-      VarRefNameOS << "::";
-      VD->getNameForDiagnostic(VarRefNameOS, Policy, true);
-    }
-    VarRefNameOS.flush();
-    if (Util::isSyclDeviceGlobalType(VD->getType())) {
-      DeviceGlobOS << "(void *)&";
-      DeviceGlobOS << VarRefName;
-    } else {
-      OS << VarRefName;
-    }
+      if (VD->isInAnonymousNamespace()) {
+        OS << TopShim;
+      } else {
+        OS << "::";
+        VD->getNameForDiagnostic(OS, Policy, true);
+      }
 
-    if (Util::isSyclDeviceGlobalType(VD->getType())) {
-      DeviceGlobOS << ", \"";
-      DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(),
-                                                          VD);
-      DeviceGlobOS << "\");\n";
-    } else {
       OS << ">() {\n";
       OS << "  return \"";
       OS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD);

From baed361e4a7d6c8f94802e0a19f2513d68e2386c Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Mon, 28 Feb 2022 15:35:57 +0300
Subject: [PATCH 10/11] Make a simple handler, align device_global

---
 clang/include/clang/Basic/Attr.td                          | 1 +
 clang/lib/Sema/SemaDeclAttr.cpp                            | 3 ---
 clang/test/CodeGenSYCL/Inputs/sycl.hpp                     | 2 +-
 clang/test/CodeGenSYCL/device_global_int_footer_header.cpp | 2 +-
 4 files changed, 3 insertions(+), 5 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 4f3b63d3558a5..5fcf47878ae73 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1267,6 +1267,7 @@ def SYCLDeviceGlobal : InheritableAttr {
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   // Only used internally by the SYCL implementation
   let Documentation = [Undocumented];
+  let SimpleHandler = 1;
 }
 
 // Marks functions which must not be vectorized via horizontal SIMT widening,
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 86e7eae9176db..d4f734d9c6c79 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -10429,9 +10429,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLSpecialClass:
     handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
     break;
-  case ParsedAttr::AT_SYCLDeviceGlobal:
-    handleSimpleAttribute<SYCLDeviceGlobalAttr>(S, D, AL);
-    break;
   case ParsedAttr::AT_SYCLDevice:
     handleSYCLDeviceAttr(S, D, AL);
     break;
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 143e8743a4822..facf6dfa2fd16 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -134,7 +134,7 @@ template <typename T>
 class [[__sycl_detail__::device_global]] device_global {
 public:
   const T &get() const noexcept { return *Data; }
-  device_global() = default;
+  device_global() {}
 
 private:
   T *Data;
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index 44ffe0f72f281..a6cd526fd6ece 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -104,7 +104,7 @@ namespace {
 struct HasVarTemplate {
   constexpr HasVarTemplate() {}
   template <typename T>
-  static constexpr device_global<T> VarTempl{};
+  static const device_global<T> VarTempl;
 };
 
 } // namespace

From 6e7b77e00da3875b171e3c6685d2af53c60bee07 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Wed, 2 Mar 2022 14:22:10 +0300
Subject: [PATCH 11/11] Remove compilation of integration footer and header

---
 .../CL/sycl/detail/defines_elementary.hpp     |  8 ---
 .../CL/sycl/detail/device_global_map.hpp      | 13 -----
 .../Inputs/CL/sycl/detail/kernel_desc.hpp     | 54 -------------------
 .../CL/sycl/detail/spec_const_integration.hpp | 15 ------
 .../device_global_int_footer_header.cpp       |  6 +--
 .../device_globals_with_spec_ids.cpp          |  3 --
 6 files changed, 1 insertion(+), 98 deletions(-)
 delete mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
 delete mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
 delete mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
 delete mode 100644 clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp

diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
deleted file mode 100644
index 00322e60ffc51..0000000000000
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp
+++ /dev/null
@@ -1,8 +0,0 @@
-#pragma once
-
-#ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
-#define __SYCL_INLINE_NAMESPACE(X) inline namespace X
-#else
-#define __SYCL_INLINE_NAMESPACE(X) namespace X
-#endif // __SYCL_DISABLE_NAMESPACE_INLINE__
-#define __SYCL_DLL_LOCAL
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
deleted file mode 100644
index a2897f48b419a..0000000000000
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/device_global_map.hpp
+++ /dev/null
@@ -1,13 +0,0 @@
-#pragma once
-
-__SYCL_INLINE_NAMESPACE(cl) {
-  namespace sycl {
-  namespace detail {
-  namespace device_global_map {
-
-  void add(void *DeviceGlobalPtr, const char *UniqueId);
-
-  } // namespace device_global_map
-  } // namespace detail
-  } // namespace sycl
-} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
deleted file mode 100644
index c4d9b998e3cb2..0000000000000
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp
+++ /dev/null
@@ -1,54 +0,0 @@
-#pragma once
-
-#include <CL/sycl/detail/defines_elementary.hpp>
-
-__SYCL_INLINE_NAMESPACE(cl) {
-  namespace sycl {
-  namespace detail {
-
-#if __cplusplus >= 201703L
-  template <auto &SpecName> const char *get_spec_constant_symbolic_ID_impl();
-  template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
-#endif
-
-#ifndef __SYCL_DEVICE_ONLY__
-#define _Bool bool
-#endif
-
-  // kernel parameter kinds
-  enum class kernel_param_kind_t {
-    kind_accessor = 0,
-    kind_std_layout = 1, // standard layout object parameters
-    kind_sampler = 2,
-    kind_pointer = 3,
-    kind_specialization_constants_buffer = 4,
-    kind_stream = 5,
-    kind_invalid = 0xf, // not a valid kernel kind
-  };
-
-  // describes a kernel parameter
-  struct kernel_param_desc_t {
-    // parameter kind
-    kernel_param_kind_t kind;
-    // kind == kind_std_layout
-    //   parameter size in bytes (includes padding for structs)
-    // kind == kind_accessor
-    //   access target; possible access targets are defined in access/access.hpp
-    int info;
-    // offset of the captured value of the parameter in the lambda or function
-    // object
-    int offset;
-  };
-
-  template <class KernelNameType> struct KernelInfo {
-    static constexpr unsigned getNumParams() { return 0; }
-    static const kernel_param_desc_t &getParamDesc(int) {
-      static kernel_param_desc_t Dummy;
-      return Dummy;
-    }
-    static constexpr const char *getName() { return ""; }
-    static constexpr bool isESIMD() { return 0; }
-  };
-  } // namespace detail
-  } // namespace sycl
-} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp b/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
deleted file mode 100644
index c3d865352ffa5..0000000000000
--- a/clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp
+++ /dev/null
@@ -1,15 +0,0 @@
-__SYCL_INLINE_NAMESPACE(cl) {
-  namespace sycl {
-  namespace detail {
-
-#if __cplusplus >= 201703L
-  // Translates SYCL 2020 `specialization_id` to a unique symbolic identifier
-  // which is used internally by the toolchain
-  template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
-    return get_spec_constant_symbolic_ID_impl<SpecName>();
-  }
-#endif
-
-  } // namespace detail
-  } // namespace sycl
-} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
index a6cd526fd6ece..d92fe72805700 100644
--- a/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
+++ b/clang/test/CodeGenSYCL/device_global_int_footer_header.cpp
@@ -2,12 +2,8 @@
 // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
 // RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
 
-// Try and compile generated integration header and footer on host.
-// RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
-
 // This test checks that integration header and footer are emitted correctly
-// for device_global variables. It also checks that emitted constructs
-// are syntactically correct.
+// for device_global variables.
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
index 674401c0d79a6..2f749b64be984 100644
--- a/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
+++ b/clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp
@@ -1,9 +1,6 @@
 // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.footer.h %s
 
-// Try and compile generated integration header and footer on host.
-// RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
-
 // This test checks that integration footer is emitted correctly if both
 // spec constants and device globals are used.