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

[tra]

LGTM.

clang/test/SemaCUDA/global-initializers-host.cu

2 ↗

(On Diff #67421)

I'd add the same comment as in the other test file explaining why both -verify is used with -emit-llvm instead of -fsyntax-only.

[jlebar]

Test tweak per review.

clang/test/SemaCUDA/global-initializers-host.cu

2 ↗

(On Diff #67421)

Good call, I don't actually need it here. :)

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

def err_config_scalar_return : Error<
   "CUDA special function 'cudaConfigureCall' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;
 def err_global_call_not_config : Error<
   "call to global function %0 not configured">;
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def err_ref_bad_target_global_initializer : Error<
+  "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
+  "function %1 in global initializer">;
 def warn_kern_is_method : Extension<
   "kernel function %0 is a member function; this may not be accepted by nvcc">,
   InGroup<CudaCompat>;
 def warn_kern_is_inline : Warning<
   "ignored 'inline' attribute on kernel function %0">,
   InGroup<CudaCompat>;
 def err_variadic_device_fn : Error<
   "CUDA device code does not support variadic functions">;

cfe/trunk/lib/Sema/SemaDecl.cpp

if (VD->isStaticLocal()) {
     }
   }
 
   // Perform check for initializers of device-side global variables.
   // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
   // 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>())) {
+    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 =
               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();
+          }
+        }
+      }
     }
   }
 
   // Grab the dllimport or dllexport attribute off of the VarDecl.
   const InheritableAttr *DLLAttr = getDLLAttr(VD);
 
   // Imported static data members cannot be defined out-of-line.
   if (const auto *IA = dyn_cast_or_null<DLLImportAttr>(DLLAttr)) {

cfe/trunk/test/SemaCUDA/global-initializers-host.cu

+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
+
+#include "Inputs/cuda.h"
+
+// Check that we get an error if we try to call a __device__ function from a
+// module initializer.
+
+struct S {
+  __device__ S() {}
+  // expected-note@-1 {{'S' declared here}}
+};
+
+S s;
+// expected-error@-1 {{reference to __device__ function 'S' in global initializer}}
+
+struct T {
+  __host__ __device__ T() {}
+};
+T t;  // No error, this is OK.
+
+struct U {
+  __host__ U() {}
+  __device__ U(int) {}
+  // expected-note@-1 {{'U' declared here}}
+};
+U u(42);
+// expected-error@-1 {{reference to __device__ function 'U' in global initializer}}
+
+__device__ int device_fn() { return 42; }
+// expected-note@-1 {{'device_fn' declared here}}
+int n = device_fn();
+// expected-error@-1 {{reference to __device__ function 'device_fn' in global initializer}}