[SYCL] Add clang support for device_global

clang/include/clang/Basic/Attr.td

@@ -1432,6 +1432,24 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
+def SYCLDeviceGlobal: InheritableAttr {
+  let Spellings = [GNU<"device_global">,
+                   CXX11<"__sycl_detail__", "device_global">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Documentation = [SYCLDeviceGlobalAttrDocs];
+  let SimpleHandler = 1;
+}
+
+def SYCLGlobalVariableAllowed : InheritableAttr {
+  let Spellings = [GNU<"global_variable_allowed">,
+                   CXX11<"__sycl_detail__", "global_variable_allowed">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Documentation = [SYCLGlobalVariableAllowedAttrDocs];
+  let SimpleHandler = 1;
+}
+
 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 SYCLDeviceGlobalAttrDocs : 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 SYCLGlobalVariableAllowedAttrDocs : 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_sycl_device_global_incorrect_scope : Error<
+  "`device_global` variables must be static or declared at namespace scope">;
+def err_not_publicly_accessible: Error<
+  "member variable %0 not publicly accessible from namespace scope">;
+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

@@ -13125,6 +13125,9 @@ class Sema final {
                                    SourceLocation BuiltinLoc,
                                    SourceLocation RParenLoc);
 
+  bool isSyclGlobalVariableAllowedType(QualType Ty);
+  bool isSyclDeviceGlobalType(QualType Ty);
+
 private:
   bool SemaBuiltinPrefetch(CallExpr *TheCall);
   bool SemaBuiltinAllocaWithAlign(CallExpr *TheCall);

clang/lib/CodeGen/CodeGenModule.cpp

@@ -2839,6 +2839,12 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
     Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
 }
 
+void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV,
+                                    const VarDecl *VD) {
+  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 +4933,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<SYCLDeviceGlobalAttr>())
+      addSYCLUniqueID(GV, D);
+  }
+
   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 VarDecl *VD);
+
   bool isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn,
                           SourceLocation Loc) const;
 

clang/lib/Sema/Sema.cpp

@@ -1621,6 +1621,18 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
   }
 }
 
+bool Sema::isSyclGlobalVariableAllowedType(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<SYCLGlobalVariableAllowedAttr>();
+  }
+  return RecTy->hasAttr<SYCLGlobalVariableAllowedAttr>();
+}
+
 namespace {
 
 /// Helper class that emits deferred diagnostic messages if an entity directly
@@ -1691,7 +1703,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) &&
+            !S.isSyclGlobalVariableAllowedType(VD->getType())) {
           S.Diag(Loc, diag::err_sycl_restrict)
               << Sema::KernelConstStaticVariable;
           return;

clang/lib/Sema/SemaDecl.cpp

@@ -7057,6 +7057,17 @@ static void copyAttrFromTypedefToDecl(Sema &S, Decl *D, const TypedefType *TT) {
     D->addAttr(Clone);
   }
 }
+bool Sema::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>();
+}
 
 NamedDecl *Sema::ActOnVariableDeclarator(
     Scope *S, Declarator &D, DeclContext *DC, TypeSourceInfo *TInfo,
@@ -7394,10 +7405,16 @@ NamedDecl *Sema::ActOnVariableDeclarator(
 
   // Static variables declared inside SYCL device code must be const or
   // constexpr
-  if (getLangOpts().SYCLIsDevice)
-    if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context))
+  if (getLangOpts().SYCLIsDevice) {
+    if (isSyclDeviceGlobalType(NewVD->getType()) &&
+        SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage()) {
+      Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope);
+    }
+    if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) &&
+        !isSyclGlobalVariableAllowedType(NewVD->getType()))
       SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict)
           << Sema::KernelNonConstStaticDataVariable;
+  }
 
   switch (D.getDeclSpec().getConstexprSpecifier()) {
   case ConstexprSpecKind::Unspecified:

clang/lib/Sema/SemaExpr.cpp

@@ -227,16 +227,18 @@ 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 SYCLGlobalVariableAllowed attribute.
       if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
           VD->getStorageClass() == SC_Static &&
-          !VD->hasAttr<SYCLGlobalVarAttr>())
+          !VD->hasAttr<SYCLGlobalVarAttr>() &&
+          !isSyclGlobalVariableAllowedType(VD->getType()))
         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 SYCLGlobalVariableAllowed attribute.
       else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
-               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
+               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>() &&
+               !isSyclGlobalVariableAllowedType(VD->getType()))
         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,29 @@ struct no_alias {
   template <bool> class instance {};
 };
 } // namespace property
+
+// device_global type decorated with attributes
+template <typename T>
+class [[__sycl_detail__::device_global]]
+[[__sycl_detail__::global_variable_allowed]] device_global {
+public :
+  const T & get() const noexcept { return *Data; }
+  device_global() {}
+  operator T&() noexcept { return *Data; }
+private:
+  T *Data;
+};
+
+// decorated with only global_variable_allowed attribute
+template <typename T>
+class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed {
+public :
+  const T & get() const noexcept { return *Data; }
+  only_global_var_allowed() {}
+  operator T&() noexcept { return *Data; }
+private:
+  T *Data;
+};
 } // namespace oneapi
 } // namespace ext
 

clang/test/CodeGenSYCL/device_global.cpp

@@ -0,0 +1,64 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s
+#include "sycl.hpp"
+
+using namespace sycl::ext::oneapi;
+using namespace cl::sycl;
+queue q;
+
+device_global<int> A;
+static device_global<int> B;
+
+struct Foo {
+  static device_global<int> C;
+};
+device_global<int> Foo::C;
+// CHECK: @A = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
+// CHECK: @_ZL1B = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
+// CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
+
+device_global<int> same_name;
+namespace NS {
+  device_global<int> same_name;
+}
+// CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]]
+// CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]]
+
+// check that we don't generate `sycl-unique-id` IR attribute if class does not use
+// [[__sycl_detail__::device_global]]
+only_global_var_allowed<int> no_device_global;
+// CHECK: @no_device_global = addrspace(1) global %"class.cl::sycl::ext::oneapi::only_global_var_allowed" zeroinitializer, align 8{{$}}
+
+void foo() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name_1>([=]() {
+      (void)A;
+      (void)B;
+      (void)Foo::C;
+      (void)same_name;
+      (void)NS::same_name;
+      (void)no_device_global;
+      });
+  });
+}
+
+namespace {
+  device_global<int> same_name;
+}
+// CHECK: @_ZN12_GLOBAL__N_19same_nameE = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ANON_NS_ATTRS:[0-9]+]]
+
+namespace {
+void bar() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name>([=]() { int A = same_name; });
+  });
+}
+
+}
+
+
+// CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
+// CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
+// CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }
+// CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" }
+// CHECK: attributes #[[SAME_NAME_NS_ATTRS]] = { "sycl-unique-id"="_ZN2NS9same_nameE" }
+// CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" }

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -64,6 +64,18 @@ namespace ext {
 namespace oneapi {
 template <typename... properties>
 class accessor_property_list {};
+
+// device_global type decorated with attributes
+template <typename T>
+struct [[__sycl_detail__::device_global]]
+[[__sycl_detail__::global_variable_allowed]] device_global {
+public:
+  const T & get() const noexcept { return *Data; }
+  device_global() {}
+  operator T&() noexcept { return *Data; }
+private:
+  T *Data;
+};
 } // namespace oneapi
 } // namespace ext