[SYCL] Add clang support for device_global

clang/include/clang/Basic/Attr.td

@@ -1432,6 +1432,22 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
+def SYCLDetailDeviceGlobal: InheritableAttr {
+  let Spellings = [GNU<"device_global">,
+                   CXX11<"__sycl_detail__", "device_global">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Documentation = [Undocumented];
+}
+
+def SYCLDetailGlobalVariableAllowed : InheritableAttr {
+  let Spellings = [GNU<"global_variable_allowed">,
+                   CXX11<"__sycl_detail__", "global_variable_allowed">];
+  let Subjects = SubjectList<[CXXRecord, GlobalStorageNonLocalVar], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Documentation = [Undocumented];
+}
+
 def SYCLIntelNoGlobalWorkOffset : InheritableAttr {
   let Spellings = [CXX11<"intel", "no_global_work_offset">];
   let Args = [ExprArgument<"Value", /*optional*/1>];

clang/include/clang/Basic/AttrDocs.td

@@ -3083,6 +3083,44 @@ function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
   }];
 }
 
+def SYCLDetailDeviceGlobalAttrDocs : Documentation {
+  let Category = DocCatType;
+  let Heading = "__sycl_detail__::device_global";
+  let Content = [{
+This attribute is part of support for SYCL device_global feature.
+[[__sycl_detail__::device_global]] attribute is used for checking restrictions
+on variable declarations using the device_global type instead of the class name.
+We do not intend to support this as a general attribute that customer code can
+use, so we have this attribute in sycl_detail namespace.
+
+.. code-block:: c++
+
+  template<typename T>
+  struct [[__sycl_detail__::device_global]] device_global {}
+
+  device_global<int> Foo;
+  }];
+}
+
+def SYCLDetailGlobalVariableAllowedAttrDocs : Documentation {
+  let Category = DocCatType;
+  let Heading = "__sycl_detail__::global_variable_allowed";
+  let Content = [{
+This attribute is part of support for SYCL device_global feature.
+[[__sycl_detail__::global_variable_allowed]] attribute is used to avoid
+diagnosing an error when variables of type device_global are referenced in
+device code. We do not intend to support this as a general attribute that
+customer code can use, therefore it is wrapped in sycl_detail namespace.
+
+.. code-block:: c++
+
+  template<typename T>
+  struct [[__sycl_detail__::device_global]] device_global {}
+
+  device_global<int> Foo;
+  }];
+}
+
 def SYCLFPGAPipeDocs : Documentation {
   let Category = DocCatStmt;
   let Heading = "pipe (read_only, write_only)";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -7099,6 +7099,17 @@ def warn_format_nonliteral : Warning<
   "format string is not a string literal">,
   InGroup<FormatNonLiteral>, DefaultIgnore;
 
+def err_non_static_member_use_not_allowed : Error<
+  "use of non-static member variable %0 is not allowed">;
+def err_not_publicly_accessible: Error<
+  "member variable %0 not publicly accessible">;
+def err_array_of_device_global_not_allowed : Error<
+  "array of device_global %0 is not allowed">;
+def err_shadow_variable_within_same_namespace: Error<
+  "shadow variable %0 not allowed withing the same enclosing namespace scope">;
+def err_namespace_name_shadows_namespace_containing_device_global : Error<
+  "not allowed: namespace name shadows %0 namespace which contains device_global">;
+
 def err_unexpected_interface : Error<
   "unexpected interface name %0: expected expression">;
 def err_ref_non_value : Error<"%0 does not refer to a value">;

clang/include/clang/Sema/Sema.h

@@ -10567,6 +10567,7 @@ class Sema final {
                                                 const SYCLUsesAspectsAttr &A);
   void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI,
                               Expr **Exprs, unsigned Size);
+
   /// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
   void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
                       bool IsPackExpansion);

clang/lib/CodeGen/CodeGenModule.cpp

@@ -2839,6 +2839,15 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
     Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
 }
 
+void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV,
+                                              const RecordDecl *RD) {
+  const auto *A = RD->getAttr<SYCLDetailDeviceGlobalAttr>();
+  assert(A && "no device_global attribute");
+  const VarDecl *VD = dyn_cast<VarDecl>(RD->getParent());
+  auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD);
+  GV->addAttribute("sycl-unique-id", builtinString);
+}
+
 bool CodeGenModule::isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn,
                                        SourceLocation Loc) const {
   const auto &NoSanitizeL = getContext().getNoSanitizeList();
@@ -4927,6 +4936,12 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   if (getLangOpts().SYCLIsDevice)
     addGlobalIntelFPGAAnnotation(D, GV);
 
+  if (getLangOpts().SYCLIsDevice) {
+    const RecordDecl *RD = D->getType()->getAsRecordDecl();
+    if (RD && RD->hasAttr<SYCLDetailDeviceGlobalAttr>())
+      addSYCLUniqueID(GV, RD);
+  }
+
   if (D->getType().isRestrictQualified()) {
     llvm::LLVMContext &Context = getLLVMContext();
 

clang/lib/CodeGen/CodeGenModule.h

@@ -1315,6 +1315,10 @@ class CodeGenModule : public CodeGenTypeCache {
   /// annotations are emitted during finalization of the LLVM code.
   void AddGlobalAnnotations(const ValueDecl *D, llvm::GlobalValue *GV);
 
+  /// Add "sycl-unique-id" llvm attribute for global variables marked with
+  /// SYCL device_global attribute
+  void addSYCLUniqueID(llvm::GlobalVariable *GV, const RecordDecl *RD);
+
   bool isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn,
                           SourceLocation Loc) const;
 

clang/lib/Sema/Sema.cpp

@@ -1691,7 +1691,8 @@ class DeferredDiagnosticsEmitter
   void visitUsedDecl(SourceLocation Loc, Decl *D) {
     if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) {
       if (auto *VD = dyn_cast<VarDecl>(D)) {
-        if (!S.checkAllowedSYCLInitializer(VD)) {
+        if (!S.checkAllowedSYCLInitializer(VD) &&
+            !VD->hasAttr<SYCLDetailGlobalVariableAllowedAttr>()) {
           S.Diag(Loc, diag::err_sycl_restrict)
               << Sema::KernelConstStaticVariable;
           return;

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4056,6 +4056,28 @@ Sema::MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A) {
   return ::new (Context) SYCLIntelLoopFuseAttr(Context, A, A.getValue());
 }
 
+static void handleSYCLDetailDeviceGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  //if (D->isATemplateDecl()) {
+  if (const auto *DeclAttr = D->getAttr<SYCLDetailDeviceGlobalAttr>()) {
+    auto *RD = dyn_cast<CXXRecordDecl>(D);
+    if (isa<FieldDecl>(RD) && !S.isUnevaluatedContext())
+      S.Diag(AL.getLoc(), diag::err_invalid_non_static_member_use) << AL;
+  }
+
+  D->addAttr(::new (S.Context) SYCLDetailDeviceGlobalAttr(S.Context, AL));
+}
+
+static void handleSYCLDetailGlobalVariableAllowedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+// Avoid diagnosing any erors here, simply accept the 
+  if (const auto *DeclAttr = D->getAttr<SYCLDetailGlobalVariableAllowedAttr>()) {
+    if (auto VD = dyn_cast<VarDecl>(D)) {
+        // avoid diagnosing error
+    }
+  }
+
+  D->addAttr(::new (S.Context) SYCLDetailGlobalVariableAllowedAttr(S.Context, AL));
+}
+
 static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) {
   // If no attribute argument is specified, set to default value '1'.
   Expr *E = A.isArgExpr(0)
@@ -10388,6 +10410,12 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
     handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLDetailDeviceGlobal:
+    handleSYCLDetailDeviceGlobalAttr(S, D, AL);
+    break;
+  case ParsedAttr::AT_SYCLDetailGlobalVariableAllowed:
+    handleSYCLDetailGlobalVariableAllowedAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
     handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL);
     break;

clang/lib/Sema/SemaExpr.cpp

@@ -227,16 +227,19 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
           (!isUnevaluatedContext() && !isConstantEvaluated());
       bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD);
       // Non-const statics are not allowed in SYCL except for ESIMD or with the
-      // SYCLGlobalVar attribute.
+      // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute.
       if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
           VD->getStorageClass() == SC_Static &&
-          !VD->hasAttr<SYCLGlobalVarAttr>())
+          (!VD->hasAttr<SYCLGlobalVarAttr>() ||
+           !VD->hasAttr<SYCLDetailGlobalVariableAllowedAttr>()))
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelNonConstStaticDataVariable;
       // Non-const globals are not allowed in SYCL except for ESIMD or with the
-      // SYCLGlobalVar attribute.
+      // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute.
       else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
-               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
+               VD->hasGlobalStorage() &&
+               (!VD->hasAttr<SYCLGlobalVarAttr>() ||
+                !VD->hasAttr<SYCLDetailGlobalVariableAllowedAttr>()))
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelGlobalVariable;
       // ESIMD globals cannot be used in a SYCL context.

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -129,6 +129,11 @@ struct no_alias {
   template <bool> class instance {};
 };
 } // namespace property
+// Global type decorated with attributes
+template <typename T>
+struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global<T>
+  device_global() = default;
+};
 } // namespace oneapi
 } // namespace ext
 

clang/test/CodeGenSYCL/device_global.cpp

@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -std=c++17 -emit-llvm %s -o - | FileCheck %s
+#include "Inputs/sycl.hpp"
+
+using namespace sycl::ext::oneapi;
+static device_global<int> Foo;
+
+device_global<int> a;           // OK
+static device_global<int> b;    // OK
+//inline device_global<int> c;    // OK
+
+struct Foo {
+  static device_global<int> d;  // OK
+};
+device_global<int> Foo::d;
+
+struct Bar {
+  device_global<int> e;         // ILLEGAL: non-static member variable not
+};                              // allowed
+
+//struct Baz {
+// private:
+//  static device_global<int> f;  // ILLEGAL: not publicly accessible from
+//};                              // namespace scope
+//device_global<int> Baz::f;
+
+//device_global<int[4]> g;        // OK
+//device_global<int> h[4];        // ILLEGAL: array of "device_global" not
+                                // allowed
+
+//device_global<int> same_name;   // OK
+//namespace foo {
+//  device_global<int> same_name; // OK
+//}
+//namespace {
+//  device_global<int> same_name; // OK
+//}
+
+//inline namespace other {
+//  device_global<int> same_name; // ILLEGAL: shadows "device_global" variable
+//}                               // with same name in enclosing namespace scope
+
+//inline namespace {
+//  namespace foo {               // ILLEGAL: namespace name shadows "::foo"
+//  }                             // namespace which contains "device_global"
+                                // variable.
+//}