[CUDA] Check initializers of instantiated template variables.

[tra]

Summary

We were already performing checks on non-template variables,
but the checks on templated ones were missing.

include/clang/Sema/Sema.h

bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                                bool ConstRHS,
                                                bool Diagnose);
 
   /// \return true if \p CD can be considered empty according to CUDA
   /// (E.2.3.1 in CUDA 7.5 Programming guide).
   bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
   bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
 
+  // \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
+  // case of error emits appropriate diagnostic and invalidates \p Var.
+  //
+  // \details CUDA allows only empty constructors as initializers for global
+  // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
+  // __shared__ variables whether they are local or not (they all are implicitly
+  // static in CUDA). One exception is that CUDA allows constant initializers
+  // for __constant__ and __device__ variables.
+  void checkAllowedCUDAInitializer(VarDecl *Var);
+
   /// Check whether NewFD is a valid overload for CUDA. Emits
   /// diagnostics and invalidates NewFD if not.
   void checkCUDATargetOverload(FunctionDecl *NewFD,
                                const LookupResult &Previous);
   /// Copies target attributes from the template TD to the function FD.
   void inheritCUDATargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD);
 
   /// \name Code completion

lib/Sema/SemaCUDA.cpp

if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
           return isEmptyCudaDestructor(Loc, RD->getDestructor());
         return true;
       }))
     return false;
 
   return true;
 }
 
+void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
+  if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
+    return;
+  const Expr *Init = VD->getInit();
+  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();
+      }
+    }
+  }
+}
+
 // With -fcuda-host-device-constexpr, an unattributed constexpr function is
 // treated as implicitly __host__ __device__, unless:
 //  * it is a variadic function (device-side variadic functions are not
 //    allowed), or
 //  * a __device__ function with this signature was already declared, in which
 //    case in which case we output an error, unless the __device__ decl is in a
 //    system header, in which case we leave the constexpr function unattributed.
 //

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) {
-    const Expr *Init = VD->getInit();
-    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();
-          }
-        }
-      }
-    }
-  }
+  if (getLangOpts().CUDA)
+    checkAllowedCUDAInitializer(VD);
 
   // 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)) {
     if (VD->isStaticDataMember() && VD->isOutOfLine() &&
         VD->isThisDeclarationADefinition()) {

lib/Sema/SemaTemplateInstantiateDecl.cpp

if (OldVar->getInit()) {
     }
 
     // We'll add an initializer to a for-range declaration later.
     if (Var->isCXXForRangeDecl() || Var->isObjCForDecl())
       return;
 
     ActOnUninitializedDecl(Var);
   }
+
+  if (getLangOpts().CUDA)
+    checkAllowedCUDAInitializer(Var);
 }
 
 /// \brief Instantiate the definition of the given variable from its
 /// template.
 ///
 /// \param PointOfInstantiation the point at which the instantiation was
 /// required. Note that this is not precisely a "point of instantiation"
 /// for the variable, but it's close.

test/SemaCUDA/device-var-init.cu

#ifdef __CUDA_ARCH__
   // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
 #endif
 }
 
 inline __host__ __device__ void hd_emitted_host_only() {
   static int x = 42; // no error on device because this is never codegen'ed there.
 }
 void call_hd_emitted_host_only() { hd_emitted_host_only(); }
+
+// Verify that we also check field initializers in instantiated structs.
+struct NontrivialInitializer {
+  __host__ __device__ NontrivialInitializer() : x(43) {}
+  int x;
+};
+
+template <typename T>
+__global__ void bar() {
+  __shared__ T bad;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+}
+
+void instantiate() {
+  bar<NontrivialInitializer><<<1, 1>>>();
+// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
+}