[sycl] [clang] Add sycl global var attribute

clang/include/clang/Basic/Attr.td

@@ -1197,6 +1197,14 @@ def SYCLDevice : InheritableAttr {
   let Documentation = [SYCLDeviceDocs];
 }
 
+def SYCLGlobalVar : InheritableAttr {
+  let Spellings = [GNU<"sycl_global_var">];
+  let Subjects = SubjectList<[GlobalVar]>;
+  let LangOpts = [SYCLIsDevice];
+  let Documentation = [SYCLGlobalVarDocs];
+  let SimpleHandler = 1;
+}
+
 def SYCLKernel : InheritableAttr {
   let Spellings = [Clang<"sycl_kernel">];
   let Subjects = SubjectList<[FunctionTmpl]>;

clang/include/clang/Basic/AttrDocs.td

@@ -3155,6 +3155,30 @@ implicitly inherit this attribute.
   }];
 }
 
+def SYCLGlobalVarDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+Normally, a SYCL kernel cannot access a global variable, but there are cases
+where it is desirable to use a global variable allocated and accessed on a SYCL device. This
+attribute is only available to a SYCL device compiler (that is, when passing
+``-fsycl-is-device``) and only applies to global variables. It affects semantic
+checks to allow use of a marked global within a SYCL kernel.
+
+.. code-block:: c++
+
+  #ifdef __SYCL_DEVICE_ONLY__
+  __attribute__((sycl_global_var)) int Var;
+  #endif
+
+  void F1(cl::sycl::handler& CGH) {
+    CGH.parallel_for_impl<class TU1>([=] () {
+      Var = 42; // device code
+    });
+  }
+
+  }];
+}
+
 def RISCVInterruptDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "interrupt (RISCV)";

clang/lib/Sema/SemaExpr.cpp

@@ -223,9 +223,10 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
       if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static)
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelNonConstStaticDataVariable;
-      // Non-const globals are allowed for SYCL explicit SIMD.
+      // Non-const globals are allowed for SYCL explicit SIMD or with the
+      // SYCLGlobalVar attribute.
       else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
-               VD->hasGlobalStorage())
+               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelGlobalVariable;
       // ESIMD globals cannot be used in a SYCL context.

clang/test/Misc/pragma-attribute-supported-attributes-list.test

@@ -152,6 +152,7 @@
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLGlobalVar (SubjectMatchRule_variable_is_global)
 // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function)

clang/test/SemaSYCL/attr-syclglobalvar.cpp

@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+
+#include "Inputs/sycl.hpp"
+
+__attribute__((sycl_global_var)) int GlobalWithAttribute;
+
+__attribute__((sycl_global_var)) extern int ExternGlobalWithAttribute;
+
+namespace NS {
+    __attribute__((sycl_global_var)) int NSGlobalWithAttribute;
+}
+
+union U {
+    int InstanceMember;
+};
+
+__attribute__((sycl_global_var)) U GlobalUnion;
+
+struct S {
+    __attribute__((sycl_global_var)) static int StaticMember;
+
+    // expected-warning@+1 {{attribute only applies to global variables}}
+    __attribute__((sycl_global_var)) int InstanceMember;
+};
+int S::StaticMember = 0;
+
+__attribute__((sycl_global_var)) S GlobalStruct;
+
+__attribute__((sycl_global_var)) static S StaticGlobal;
+
+// expected-error@+1 {{attribute takes no arguments}}
+__attribute__((sycl_global_var(42))) int GlobalWithAttributeArg;
+
+int GlobalNoAttribute;
+
+// expected-warning@+1 {{attribute only applies to global variables}}
+__attribute__((sycl_global_var)) void F() {
+    __attribute__((sycl_global_var)) static int StaticLocalVar;
+
+    // expected-warning@+1 {{attribute only applies to global variables}}
+    __attribute__((sycl_global_var)) int Local;
+
+    cl::sycl::kernel_single_task<class kernel_name>([=] () {
+        (void)GlobalWithAttribute;
+        (void)ExternGlobalWithAttribute;
+        (void)NS::NSGlobalWithAttribute;
+        (void)GlobalUnion.InstanceMember;
+        (void)S::StaticMember;
+        (void)GlobalStruct.InstanceMember;
+        (void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
+        (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
+        (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}}
+    });
+}