[SYCL] Add the remaining diagnostics to device_global implementation

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -7219,9 +7219,11 @@ def warn_format_nonliteral : Warning<
   InGroup<FormatNonLiteral>, DefaultIgnore;
 
 def err_sycl_device_global_incorrect_scope : Error<
-  "'device_global' variables must be static or declared at namespace scope">;
+  "'device_global' variable must be a static data member or declared in global or namespace scope">;
 def err_sycl_device_global_not_publicly_accessible: Error<
-  "'device_global' member variable %0 is not publicly accessible from namespace scope">;
+  "'device_global' member variable %0 should be publicly accessible from namespace scope">;
+def err_sycl_device_global_array : Error<
+  "'device_global' array is not allowed">;
 
 def err_unexpected_interface : Error<
   "unexpected interface name %0: expected expression">;

clang/lib/Sema/SemaDecl.cpp

@@ -7678,13 +7678,32 @@ NamedDecl *Sema::ActOnVariableDeclarator(
       NewVD->setTSCSpec(TSCS);
   }
 
-  // Global variables with types decorated with device_global attribute must be
-  // static if they are declared in SYCL device code.
   if (getLangOpts().SYCLIsDevice) {
-    if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() &&
-        isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
-            NewVD->getType()))
-      Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope);
+    // device_global array is not allowed.
+    if (const ArrayType *AT = getASTContext().getAsArrayType(NewVD->getType()))
+      if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+              AT->getElementType()))
+        Diag(NewVD->getLocation(), diag::err_sycl_device_global_array);
+
+    // Global variables with types decorated with device_global attribute must
+    // be static if they are declared in SYCL device code.
+    if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            NewVD->getType())) {
+      if (SCSpec == DeclSpec::SCS_static) {
+        const DeclContext *DC = NewVD->getDeclContext();
+        while (!DC->isTranslationUnit()) {
+          if (isa<FunctionDecl>(DC)) {
+            Diag(D.getIdentifierLoc(),
+                 diag::err_sycl_device_global_incorrect_scope);
+            break;
+          }
+          DC = DC->getParent();
+        }
+      } else if (!NewVD->hasGlobalStorage()) {
+        Diag(D.getIdentifierLoc(),
+             diag::err_sycl_device_global_incorrect_scope);
+      }
+    }
 
     // Static variables declared inside SYCL device code must be const or
     // constexpr unless their types are decorated with global_variable_allowed

clang/lib/Sema/SemaDeclCXX.cpp

@@ -3598,10 +3598,24 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D,
   if (getLangOpts().SYCLIsDevice) {
     if (auto Value = dyn_cast<ValueDecl>(Member)) {
       if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
-              Value->getType()) &&
-          Value->getAccess() != AS_public) {
-        Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
-            << Value;
+              Value->getType())) {
+        if (Value->getAccess() == AS_private ||
+            Value->getAccess() == AS_protected) {
+          Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
+              << Value;
+        }
+        const DeclContext *DC = Member->getDeclContext();
+        while (!DC->isTranslationUnit()) {
+          if (auto Decl = dyn_cast<NamedDecl>(DC)) {
+            if (Decl->getAccess() == AS_private ||
+                Decl->getAccess() == AS_protected) {
+              Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
+                  << Value;
+              break;
+            }
+          }
+          DC = DC->getParent();
+        }
       }
     }
   }

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -1614,14 +1614,38 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
   // Only add this if we aren't instantiating a variable template.  We'll end up
   // adding the VarTemplateSpecializationDecl later.
   if (!InstantiatingVarTemplate) {
-    SemaRef.addSyclVarDecl(Var);
+    if (SemaRef.getLangOpts().SYCLIsDevice &&
+        SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            Var->getType())) {
+      if (!Var->hasGlobalStorage())
+        SemaRef.Diag(D->getLocation(),
+                     diag::err_sycl_device_global_incorrect_scope);
+
+      if (Var->getAccess() == AS_private || Var->getAccess() == AS_protected)
+        SemaRef.Diag(D->getLocation(),
+                     diag::err_sycl_device_global_not_publicly_accessible)
+            << Var;
+
+      if (Var->isStaticLocal()) {
+        const DeclContext *DC = Var->getDeclContext();
+        while (!DC->isTranslationUnit()) {
+          if (isa<FunctionDecl>(DC)) {
+            SemaRef.Diag(D->getLocation(),
+                         diag::err_sycl_device_global_incorrect_scope);
+            break;
+          }
+          DC = DC->getParent();
+        }
+      }
+    }
     if (const auto *SYCLDevice = Var->getAttr<SYCLDeviceAttr>()) {
       if (!SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
               Var->getType()))
         SemaRef.Diag(SYCLDevice->getLoc(),
                      diag::err_sycl_attribute_not_device_global)
             << SYCLDevice;
     }
+    SemaRef.addSyclVarDecl(Var);
   }
   return Var;
 }
@@ -1711,6 +1735,17 @@ Decl *TemplateDeclInstantiator::VisitFieldDecl(FieldDecl *D) {
 
   Field->setImplicit(D->isImplicit());
   Field->setAccess(D->getAccess());
+  // Static members are not processed here, so error out if we have a device
+  // global without checking access modifier.
+  if (SemaRef.getLangOpts().SYCLIsDevice) {
+    if (SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            Field->getType())) {
+      SemaRef.Diag(D->getLocation(),
+                   diag::err_sycl_device_global_incorrect_scope);
+      Field->setInvalidDecl();
+      return nullptr;
+    }
+  }
   Owner->addDecl(Field);
 
   return Field;

clang/test/SemaSYCL/device_global.cpp

@@ -1,9 +1,10 @@
 // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s
 #include "Inputs/sycl.hpp"
 
-// Test cases below check for valid usage of device_global and
-// global_variable_allowed attributes, and that they are being correctly
-// generated in the AST.
+// Diagnostic tests for device_global and global_variable_allowed attribute.
+
+// Test that there are no errors when variables of type device_global are
+// decorated with global_variable_allowed attribute appropriately.
 using namespace sycl::ext::oneapi;
 
 device_global<int> glob;                  // OK
@@ -18,19 +19,95 @@ device_global<char> Foo::d;
 
 struct Baz {
 private:
-  // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}}
+  // expected-error@+1{{'device_global' member variable 'f' should be publicly accessible from namespace scope}}
   static device_global<int> f;
+
+protected:
+  // expected-error@+1{{'device_global' member variable 'g' should be publicly accessible from namespace scope}}
+  static device_global<int> g;
 };
+
 device_global<int> Baz::f;
 
 device_global<int[4]> not_array; // OK
 
+// expected-error@+1{{'device_global' array is not allowed}}
+device_global<int> array[4];
+
 device_global<int> same_name; // OK
+
 namespace foo {
 device_global<int> same_name; // OK
 }
-namespace {
-device_global<int> same_name; // OK
+
+struct BBar {
+private:
+  struct BarInsider {
+    // expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
+    static device_global<float> c;
+  };
+
+protected:
+  struct BarInsiderProtected {
+    // expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
+    static device_global<float> c;
+  };
+};
+
+struct ABar {
+  void method() {
+    // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+    static device_global<float> c;
+  }
+  struct BarInsider {
+    static device_global<float> c;
+    void method() {
+      // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+      static device_global<float> c;
+    }
+  };
+};
+
+template <typename T> void fooBar() {
+  // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  static device_global<T> c;
+  // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<T> d;
+}
+
+template <typename T> struct TS {
+private:
+  // expected-error@+1 2{{'device_global' member variable 'a' should be publicly accessible from namespace scope}}
+  static device_global<T> a;
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<T> b;
+  // expected-error@+2{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<int> c;
+
+public:
+  static device_global<T> d;
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<T> e;
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<int> f;
+
+protected:
+  // expected-error@+1 2{{'device_global' member variable 'g' should be publicly accessible from namespace scope}}
+  static device_global<T> g;
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<T> h;
+  // expected-error@+2{{'device_global' member variable 'i' should be publicly accessible from namespace scope}}
+  // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
+  device_global<int> i;
+};
+
+// expected-note@+1{{in instantiation of template class 'TS<int>' requested here}}
+TS<int> AAAA;
+
+//expected-note@+2{{in instantiation of template class 'TS<char>' requested here}}
+template <typename T> void templFoo () {
+  TS<T> Var;
 }
 
 // expected-error@+2{{'device_global' attribute only applies to classes}}
@@ -44,6 +121,12 @@ device_global<int> same_name; // OK
 union [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] a_union;
 
 int main() {
+  // expected-note@+1{{in instantiation of function template specialization 'templFoo<char>' requested here}}
+  templFoo<char>();
+
+  // expected-note@+1{{in instantiation of function template specialization 'fooBar<int>' requested here}}
+  fooBar<int>();
+
   sycl::kernel_single_task<class KernelName1>([=]() {
     (void)glob;
     (void)static_glob;
@@ -53,11 +136,7 @@ int main() {
   });
 
   sycl::kernel_single_task<class KernelName2>([]() {
-    // expected-error@+1{{'device_global' variables must be static or declared at namespace scope}}
+    // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
     device_global<int> non_static;
-
-    // expect no error on non_const_static declaration if decorated with
-    // [[__sycl_detail__::global_variable_allowed]]
-    static device_global<int> non_const_static;
   });
 }