[CUDA] Reject calls to __device__ functions from host variable global initializers.

Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D23335

llvm-svn: 278196

1 files changed, 47 insertions, 28 deletions

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index a18ef0c4cf3..e850c3c3847 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -10728,36 +10728,55 @@ Sema::FinalizeDeclaration(Decl *ThisDecl) {
   // 7.5). We must also apply the same checks to all __shared__
   // variables whether they are local or not. CUDA also allows
   // constant initializers for __constant__ and __device__ variables.
-  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+  if (getLangOpts().CUDA) {
     const Expr *Init = VD->getInit();
-    if (Init && VD->hasGlobalStorage() &&
-        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
-         VD->hasAttr<CUDASharedAttr>())) {
-      assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
-      bool AllowedInit = false;
-      if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
-        AllowedInit =
-            isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
-      // We'll allow constant initializers even if it's a non-empty
-      // constructor according to CUDA rules. This deviates from NVCC,
-      // but allows us to handle things like constexpr constructors.
-      if (!AllowedInit &&
-          (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
-        AllowedInit = VD->getInit()->isConstantInitializer(
-            Context, VD->getType()->isReferenceType());
-
-      // Also make sure that destructor, if there is one, is empty.
-      if (AllowedInit)
-        if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+    if (Init && VD->hasGlobalStorage()) {
+      if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
+          VD->hasAttr<CUDASharedAttr>()) {
+        assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
+        bool AllowedInit = false;
+        if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
           AllowedInit =
-              isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
-
-      if (!AllowedInit) {
-        Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
-                                    ? diag::err_shared_var_init
-                                    : diag::err_dynamic_var_init)
-            << Init->getSourceRange();
-        VD->setInvalidDecl();
+              isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+        // We'll allow constant initializers even if it's a non-empty
+        // constructor according to CUDA rules. This deviates from NVCC,
+        // but allows us to handle things like constexpr constructors.
+        if (!AllowedInit &&
+            (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+          AllowedInit = VD->getInit()->isConstantInitializer(
+              Context, VD->getType()->isReferenceType());
+
+        // Also make sure that destructor, if there is one, is empty.
+        if (AllowedInit)
+          if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+            AllowedInit =
+                isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+
+        if (!AllowedInit) {
+          Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
+                                      ? diag::err_shared_var_init
+                                      : diag::err_dynamic_var_init)
+              << Init->getSourceRange();
+          VD->setInvalidDecl();
+        }
+      } else {
+        // This is a host-side global variable.  Check that the initializer is
+        // callable from the host side.
+        const FunctionDecl *InitFn = nullptr;
+        if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
+          InitFn = CE->getConstructor();
+        } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
+          InitFn = CE->getDirectCallee();
+        }
+        if (InitFn) {
+          CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
+          if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
+            Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
+                << InitFnTarget << InitFn;
+            Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
+            VD->setInvalidDecl();
+          }
+        }
       }
     }
   }